7.4 Overhead, latency hiding וחוק Little הרצאה

בשלושת השיעורים הקודמים של הפרק בנינו את התמונה הגדולה של ניתוח ביצועים. בשיעור 7.1 הגדרנו צוואר-בקבוק (bottleneck) ואת מודל ה-Roofline, בשיעור 7.2 הפרדנו בין kernel חסום-חישוב (compute-bound) לבין kernel חסום-זיכרון (memory-bound), ובשיעור 7.3 מדדנו את ה-arithmetic intensity (עצימות אריתמטית) שקובעת באיזה צד של נקודת-הרכס אנחנו יושבים. אבל שמנו לב לפרט מטריד: kernels אמיתיים כמעט אף פעם לא נוגעים בגג. הם יושבים מתחתיו, ומודל ה-Roofline - שמצייר רק רוחבי-פס ותפוקות - לא מסביר למה. הסיבה היא הקטגוריה השלישית של העבודה, זו שלא מופיעה על ה-Roofline כלל: תקורה (overhead). בשיעור הזה נשלים את המשולש. נגדיר תקורה כזמן שבו ה-GPU יושב בטל וממתין לעבודה, נבדיל אותה חדות מצוואר-בקבוק (שבו הוא עובד במלוא הכוח), ונכיר את שתי מקורות התקורה העיקריים - תקורת המארח (host, ה-CPU) ותקורת התקשורת. משם נעבור לצד השני של המטבע: הסתרת latency (latency hiding), המנגנון שבזכותו ה-GPU שומר על תפוקה גבוהה למרות ש-latency של פעולה בודדת ארוך מאוד, ולבסוף לחוק Little (Little's Law) - הנוסחה הכמותית שאומרת לנו בדיוק כמה מקביליות צריך כדי להסתיר את ה-latency הזה. חוק Little הוא היסוד התיאורטי ל-occupancy (תפוסה), שנקדיש לו את פרק 8 כולו.

התקורה - הצוואר השלישי - overhead¶

ה-glossary מגדיר overhead במשפט אחד חד: הoverhead היא הזמן שמושקע בלי שנעשית עבודה שימושית. זו ההגדרה, וכל השאר נגזר ממנה. שימו לב להבחנה שהיא יוצרת מול שני הצווארים שכבר הכרנו. כש-kernel compute-bound, ה-GPU עובד במלוא המהירות - יחידות ה-CUDA Core וה-Tensor Core רוויות, פשוט אין להן יותר מה לתת. כש-kernel memory-bound, מערכת הזיכרון רוויה - ה-HBM מזרים בייטים בקצב המרבי שלו. בשני המקרים המשאב המגביל עובד קשה. הoverhead היא ההפך הגמור: ה-GPU יושב בטל, ממתין לקבל עבודה או ממתין לתלות שתיפתר. לא עבודה איטית - היעדר עבודה.

זו בדיוק הסיבה שoverhead איננה מופיעה על ה-Roofline. ה-Roofline מציב שני גגות - רוחב-פס אריתמטי ורוחב-פס של הזיכרון - שניהם קצבים (throughput). הoverhead איננה קצב; היא latency של המתנה. שני צירי ה-Roofline (ops/byte ו-ops/s) פשוט לא יכולים לבטא אותה. לכן kernel חסום-תקורה (overhead-bound) יושב עמוק מתחת לשני הגגות: הוא לא נכשל כי החומרה איטית, אלא כי החומרה מבזבזת זמן בהמתנה. הדוגמה הקנונית של ה-glossary לעומס overhead-bound היא פעולות על מערכים קטנים - kernel שרץ מיקרו-שניות בודדות על הdevice, אבל עלה להost מיקרו-שניות בודדות רק כדי לdispatch אותו.

נחלק את העבודה של כל kernel לשלוש קטגוריות (הפירוק המיוחס להרצאה של Horace He):

   All kernel work breaks down into:

   ┌──────────────┬──────────────────────────────────────────┐
   │ Compute      │ floating-point operations on CUDA / Tensor Cores │ -> compute-bound
   ├──────────────┼──────────────────────────────────────────┤
   │ Memory       │ moving data through the memory hierarchy   │ -> memory-bound
   ├──────────────┼──────────────────────────────────────────┤
   │ Overhead     │ "everything else" - idle time, waiting, bookkeeping │ -> overhead-bound
   └──────────────┴──────────────────────────────────────────┘

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

תקורת המארח - host overhead¶

מקור הoverhead הראשון והנפוץ ביותר הוא host overhead: הזמן שה-CPU מבזבז כדי לגרום ל-GPU לעבוד. ה-glossary נותן את המספר המכונן: בערך 10 מיקרו-שניות של זמן CPU לכל launch kernel. כשאתם כותבים myKernel<<<grid, block>>>(args), ה-<<< >>> הוא סוכר תחבירי שמתורגם לרצף קריאות אל ה-Runtime API (בעיקר cudaLaunchKernel, כפי שראינו בפרק 5). הרצף הזה עושה עבודת CPU אמיתית: אריזת הארגומנטים למערך, העתקת תצורת הlaunch (grid, block, shared memory, stream), ופנייה לדרייבר שכותב פקודה לתור החומרה. כל זה עולה בסדר גודל של יחידות מיקרו-שניות - גם אם ה-kernel עצמו רץ על הdevice במיקרו-שנייה אחת.

וכשמעורב framework כמו PyTorch או TensorFlow, התמונה מחמירה: לפני הlaunch עצמה, ה-framework מבצע קבלת החלטות - איזה kernel לבחור, איזה dispatch לגזרת הטיפוס והצורה, בדיקות אוטו-גראד - וזה מוסיף מיקרו-שניות רבות נוספות על גבי ה-10 של ה-API. הoverhead של השכבה הפייתונית לבדה עלולה להיות גדולה מזמן הlaunch של CUDA.

כל עוד ה-kernels גדולים - רצים מילישניות שלמות - הoverhead זניחה. הבעיה מתלקחת כשה-kernels קטנים ורבים. נעשה חשבון על צעד decode יחיד של מודל שפה בסגנון Transformer, בן 32 שכבות:

   per layer, roughly: QKV projection, RoPE, attention, output
   projection, two GEMMs for the MLP, two normalizations, and more
   ->  roughly 15 kernels per layer

   32 layers x ~15 kernels   ≈   480 launches per token
   480 launches x ~5 microseconds ≈   2.4 milliseconds of CPU overhead alone

אם יעד ה-latency הוא 10 מילישניות לטוקן, אז 2.4 מילישניות נשרפים בoverhead לפני שה-GPU חישב ביט אחד - כמעט רבע מהתקציב. וגרוע מכך: אם ה-kernels קצרים מהזמן שלוקח להost לdispatch את הבא בתור, ה-GPU מורעב (starved) - מסיים את ה-kernel הנוכחי ויושב בטל בזמן שה-CPU עדיין טורח על הlaunch הבאה. זהו בדיוק המצב הoverhead-bound: לא חישוב איטי ולא זיכרון איטי, אלא host שלא מספיק להזין את הdevice.

כיצד מקטינים host overhead. האסטרטגיה הכללית אחת: לשלם את מחיר הlaunch פחות פעמים, על עבודה גדולה יותר. ארבע הגרסאות המעשיות:

גרפים - CUDA Graphs. לוכדים DAG שלם של launches ומגישים אותו בקריאת host אחת שאפשר להריץ מחדש (replay) בזול. זו התרופה הישירה, והקדשנו לה את שיעור 5.5. במקום 480 launches לטוקן - cudaGraphLaunch אחד.
מיזוג - fusion. מאחדים כמה פעולות ל-kernel אחד (למשל bias + activation + dropout ל-kernel יחיד). פחות kernels = פחות launches = פחות overhead, וגם פחות תעבורת זיכרון (שיעור 7.3).
קיבוץ - batching. משגרים kernel אחד שמעבד batch גדול במקום הרבה kernels קטנים. אותו מספר launches פורש על יותר עבודה - הoverhead לlaunch נשארת אבל חלקה היחסי בזמן צונח.
פחות launches, גדולות יותר. העיקרון שמאחד את כולן: כל טריק שמוריד את מספר הlaunches או מגדיל את העבודה לכל launch מקטין את חלק הoverhead.

תקורת תקשורת - communication overhead¶

מקור הoverhead השני הוא תקשורת: זמן ההשהיה בהזזת נתונים בין CPU ל-GPU, או בין GPU ל-GPU. כשמעתיקים נתונים על-פני ה-PCIe בין הhost להdevice, או על-פני NVLink בין שני GPUs, אנחנו מוציאים זמן על תנועת נתונים, לא על חישוב. ה-glossary נותן מסגור מדויק ורב-עוצמה לזה: כשרוחב-הפס של התקשורת מגביל את הביצועים, יש להתייחס לזה כאל צורה של חסימת-זיכרון (memory-boundedness) על-פני מערכת מבוזרת. במילים אחרות, אותו ניתוח Roofline שלמדנו למערכת בודדת חל גם על אשכול (cluster) שלם: ה"זיכרון" הוא הצומת המרוחק, ורוחב-הפס של הרשת (NVLink, InfiniBand) הוא הגג המשופע.

   the "bandwidth" hierarchy - from fastest to slowest:

   register file  <-> SM         : ~tens of TB/s     (inside the SM)
   HBM3           <-> SM         : ~3.35 TB/s        (H100, global memory)
   NVLink         GPU <-> GPU    : ~hundreds of GB/s  (within a node)
   PCIe Gen5      CPU <-> GPU    : ~64 GB/s           (host-device)
   InfiniBand     node <-> node  : ~tens of GB/s      (between machines)

ככל שיורדים בהיררכיה רוחב-הפס צונח בסדרי גודל. לכן העברה של tensor על-פני PCIe או על-פני הרשת עלולה בקלות להפוך לצוואר-הבקבוק של אימון או inference מבוזר - וזו הסיבה שספריות כמו NCCL מתאמצות כל-כך לחפוף תקשורת עם חישוב. התרופות דומות לhost overhead בעיקרון: פחות העברות, גדולות יותר (איחוד הודעות קטנות לגדולה), חפיפה של תקשורת עם חישוב (streams, שיעור 5.3), והשארת נתונים על הdevice ככל האפשר במקום הלוך-ושוב אל הhost.

הסתרת latency - latency hiding¶

עד כאן דיברנו על overhead שנוצרת מחוץ ל-GPU (הhost, הרשת). אבל יש latency גם בתוך ה-GPU: פקודת load מהזיכרון הגלובלי לוקחת מאות מחזורי-שעון, ואפילו פקודות אריתמטיות לוקחות כמה מחזורים. אם ה-GPU היה עוצר ומחכה לכל load, הוא היה בטל רוב הזמן. הפתרון הוא הסתרת latency: אסטרטגיה למסך פעולות ארוכות-latency על-ידי הרצת רבות מהן במקביל. זהו הרעיון המכונן של ארכיטקטורת ה-GPU, ובו נגענו כבר בשיעור 1.4 כשהכרנו את ה-warp scheduler.

התובנה המרכזית, שה-glossary מנסח בחדות: warps, לא threads, הם יחידת ההנפקה של פקודות. ("Threads are not the unit of instruction issuance, warps are.") ה-warp scheduler שומר על מאגר גדול של warps תושבים (resident) על ה-SM, ובכל מחזור בוחר warp אחד כשיר (eligible) ומנפיק לו פקודה. כשה-warp הזה נתקע (stalled) בהמתנה ל-load, המתזמן פשוט מנפיק מ-warp כשיר אחר. מיתוג ההקשר הזה כמעט חינמי - מצב האוגרים של כל warp נשאר תושב בקובץ-האוגרים של ה-SM, ולכן החלפת warp עולה בערך מחזור-שעון אחד (זה בדיוק המספר שראינו ל-H100 בשיעור 1.2). בזמן שה-load של warp אחד "בטיסה", המתזמן ממלא את החורים בעבודה של warps אחרים. ה-latency של ה-load הבודד נשאר ארוך - אבל הוא מוסתר מאחורי עבודה מועילה.

יש עוד ממד להסתרה: warps שונים יכולים לתפוס pipes שונים בו-זמנית. ב-H100 יש בכל SM ארבעה warp schedulers, וכל אחד שולט על pipelines שונות: Tensor Core, CUDA Core (FP32/INT32), ו-LSU (load/store unit). בזמן ש-warp אחד מריץ כפל-מטריצות על ה-Tensor Cores, warp שני מריץ אריתמטיקה על ה-CUDA Cores, ו-warp שלישי מושך נתונים דרך ה-LSU - כל השלושה חופפים ומסתירים זה את ה-latency של זה. זו הקבלה ל-SMT (Simultaneous Multi-Threading) של מעבדים, רק בקנה-מידה גדול בהרבה: במקום 2 threads לליבה, עשרות warps ל-SM.

נראה זאת על SASS קונקרטי - שרשרת תלויה של פקודות (כל אחת תלויה בתוצאה של קודמתה), עם ה-latency האופייני של כל פקודה:

LDG.E.SYS R1, [R0]        // load from global memory, ~400 cycles
IMUL      R2, R1, 0xBEEF  // integer multiply, ~6 cycles   (depends on R1)
IADD      R4, R2, 0xAFFE  // integer add, ~4 cycles   (depends on R2)
IMUL      R6, R4, 0x1337  // integer multiply, ~6 cycles   (depends on R4)

מכיוון שכל פקודה תלויה בקודמתה, warp בודד חייב לעצור לאורך כל שרשרת ה-latency: 400 + 6 + 4 + 6 = 416 מחזורים שבהם הוא לא יכול להתקדם. אם רק ה-warp הזה רץ על ה-SM, יחידות הביצוע יושבות בטל 416 מחזורים - בזבוז מוחלט. הדרך היחידה למלא את החור היא warps רבים ובלתי-תלויים: בעוד ה-warp הראשון תקוע על ה-load, warps אחרים מנפיקים את הפקודות שלהם. כמה warps צריך בדיוק? לזה בדיוק נועד חוק Little.

חוק Little - Little's Law¶

חוק Little הוא, בלשון ה-glossary, "החשוב שבחוקים הבסיסיים" של ניתוח מערכות כמותי. הוא קובע את כמות המקביליות הדרושה כדי להסתיר latency במלואו באמצעות throughput. הנוסחה פשוטה להפליא:

   concurrency (ops)  =  latency (s)  x  throughput (ops/s)

שלושת המרכיבים:

מקביליות - concurrency - מספר הפעולות שחייבות להיות "בטיסה" (in flight) בו-זמנית.
השהיה - latency - זמן ההשהיה של פעולה בודדת, בשניות (או במחזורים).
תפוקה - throughput - מספר הפעולות שמושלמות בשנייה (הקצב שאתם רוצים לקיים).

זו בדיוק צורת ההנדסה-מערכתית של תוצאת התורים הקלאסית L = λW: ה-concurrency הוא L (מספר הפריטים במערכת), ה-throughput הוא λ (קצב ההגעה), וה-latency הוא W (זמן השהייה). האינטואיציה: אם כל פעולה שוהה W שניות, ואתם רוצים להשלים λ פעולות בשנייה, אז בכל רגע נתון יש λ·W פעולות "תלויות באוויר" בו-זמנית. אם אין לכם מספיק פעולות בטיסה - הpipe מתרוקן וה-throughput צונח מתחת ל-λ.

נציב את המספרים של ה-glossary. נניח שאנחנו רוצים לקיים throughput של פקודה אחת למחזור, וה-latency של גישה לזיכרון הוא 400 מחזורים:

   concurrency = latency x throughput
               = 400 cycles x 1 operation/cycle
               = 400 operations in flight

צריך 400 פעולות זיכרון בטיסה בו-זמנית כדי לשמור על יחידת הביצוע עסוקה כל מחזור. עכשיו נניח שהחומרה מהירה יותר ומסוגלת ל-10 פקודות למחזור (אותו latency של 400):

   concurrency = 400 cycles x 10 operations/cycle
               = 4,000 operations in flight

זו התובנה החדה של חוק Little: throughput גבוה יותר דורש מקביליות גבוהה יותר, באופן פרופורציוני, כדי להישאר מוסתר-latency. ככל שהחומרה מהירה יותר, כך צריך יותר עבודה מקבילה "באוויר" רק כדי להזין אותה. זה מסביר מדוע GPUs, שהם מכונות throughput קיצוניות, זקוקים לכל-כך הרבה threads: לא כי כל thread מהיר, אלא כי צריך המון מהם בטיסה כדי להסתיר את ה-latency ולשמור על הקצב.

מ-latency ל-מספר warps¶

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

   warps to hide  ≈  latency (cycles) / warp width
                 =  416 / 32
                 ≈  13 warps

צריך בערך 13 warps בלתי-תלויים שרצים במקביל כדי להסתיר לגמרי את שרשרת ה-416 מחזורים. הרעיון: כשה-warp הראשון תקוע על ה-load שלו, 12 ה-warps האחרים מנפיקים את הפקודות שלהם ומכסים את החור, כך שבכל מחזור לפחות warp אחד כשיר להנפיק. שימו לב שהמספר הזה עקבי עם דוגמת חוק Little: 400 פעולות-thread בטיסה חלקי 32 threads ל-warp נותן ~12.5 ≈ 13 warps. שתי הזוויות - הפעולות בטיסה ומספר ה-warps - הן אותו חשבון בדיוק, רק ביחידות שונות.

והנה החיבור ל-occupancy, שנקדיש לו את פרק 8: 13 warps מתוך המקסימום של 64 warps ל-SM ב-H100 הם occupancy של בערך 13/64 ≈ 20%. כלומר, כבר ב-occupancy צנוע של כ-20% אנחנו מסתירים במלואה שרשרת load אופיינית. זו נקודה קריטית שנחזור אליה: לא תמיד צריך occupancy גבוה - צריך רק מספיק כדי להסתיר את ה-latency שבפועל. מעבר לזה, occupancy נוסף לא קונה כלום (ולפעמים אף מזיק, כי הוא מקטין את מספר האוגרים ל-thread). את הדקות הזו נפרוט בשיעור 8.1.

מלכודת נפוצה: החישוב latency / 32 מניח ששרשרת התלות מכריחה כל warp להחזיק פקודה אחת בטיסה בכל רגע. אם הקוד שלכם דווקא חושף מקביליות ברמת הפקודה (ILP) - למשל כמה loads בלתי-תלויים ברצף באותו warp - אז warp בודד יכול להחזיק כמה פעולות בטיסה, וצריך פחות warps כדי להגיע לאותו concurrency. ILP ו-occupancy הם שני מנופים חלופיים לאותו יעד: להחזיק מספיק פעולות בטיסה.

התוצאה של Volkov - זיכרון מול אריתמטיקה¶

עכשיו לתוצאה הכי לא-אינטואיטיבית, וגם היפה ביותר, בכל השיעור. ה-glossary מפנה לסעיף 4.3 בתזה של Vasily Volkov, שמדד ניסויית כמה warps צריך כדי להסתיר latency. הציפייה הטבעית: latency של זיכרון ארוך בהרבה מ-latency אריתמטי (מאות מחזורים מול יחידות בודדות), אז ודאי שצריך הרבה יותר warps כדי להסתיר זיכרון. התוצאה הפתיעה:

   to hide pure arithmetic latency    :  ~24 warps
   to hide pure memory latency        :  ~30 warps

כמעט אותו מספר. איך ייתכן, אם ה-latency של הזיכרון גדול פי כמה? התשובה היא בדיוק חוק Little: concurrency = latency x throughput, ושני הגורמים מתקזזים.

   arithmetic:  low latency    x  high throughput  =  moderate concurrency
   memory:      high latency   x  low throughput   =  moderate concurrency
                    │                   │
                    └── several times larger   └── several times smaller  ->  the product is similar

לאריתמטיקה יש latency נמוך אבל throughput גבוה (ה-CUDA Cores/Tensor Cores מסוגלים להמון פעולות למחזור, כפי שראינו ב-arithmetic bandwidth). לזיכרון יש latency גבוה אבל throughput נמוך (רוחב-הפס של ה-HBM נמוך בהרבה מרוחב-הפס האריתמטי - זה בדיוק ה"קיר הזיכרון" מ-7.1, והיחס Tensor Core:CUDA Core של בערך 100:1). בחוק Little, ה-concurrency הוא המכפלה של השניים. הזיכרון מפצה את ה-latency הגדול שלו ב-throughput קטן, והאריתמטיקה מפצה את ה-throughput הגדול שלה ב-latency קטן - והמכפלות יוצאות דומות. לכן מספר ה-warps הדרוש להסתרה דומה.

זו לא מקריות - זו תכונה שימושית של מכונת latency-hiding. מכיוון שאותו מספר warps מסתיר גם latency אריתמטי וגם latency של זיכרון, אותה תצורת occupancy מסתירה עומס מעורב של אריתמטיקה וזיכרון - שזה בדיוק מה שרוב ה-kernels האמיתיים עושים. ה-GPU מתוכנן כך שהמאזן הזה יעבוד עבורכם.

אבחון - overhead-bound מול compute-bound מול memory-bound¶

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

מאפיין	compute-bound	memory-bound	overhead-bound
מה המשאב המגביל	רוחב-פס אריתמטי	רוחב-פס הזיכרון	latency / המתנה
מצב יחידות הביצוע	רוויות, עובדות מלא	ממתינות לזיכרון	בטלות, אין עבודה
מופיע על ה-Roofline	כן (מתחת לגג האופקי)	כן (מתחת לגג המשופע)	לא (Roofline לא רואה latency)
מדד ncu מרכזי	Compute (SM) Throughput גבוה	Memory Throughput גבוה	שניהם נמוכים
מה מזיז את המחט	Tensor Cores, precision נמוך	coalescing, reuse, קיטול	גרפים, מיזוג, קיבוץ, occupancy
דוגמה	GEMM גדול	כפל וקטור-וקטור	פעולות על מערכים קטנים

השורה שמסגירה overhead-boundedness: גם ה-Compute Throughput וגם ה-Memory Throughput נמוכים בו-זמנית. אם שני המדדים למטה, החומרה לא עובדת קשה על שום דבר - היא בטלה. זה בדיוק ההפך משני הצווארים האחרים, שבהם לפחות משאב אחד רווי. בדיוק כפי שראינו, ה-Roofline לא יעזור לכם כאן - הוא מראה רק כמה קרובים לגג, וכשאתם overhead-bound אתם רחוקים משני הגגות בגלל סיבה שהוא לא מודד.

נחבר את זה למספרים הקונקרטיים של ה-H100 ולחוק Little. ל-H100 יש 132 SMs, ובכל SM ארבעה warp schedulers שכל אחד מנפיק עד פקודה אחת למחזור:

   maximum issue rate, H100 @ 1.98 GHz:
   132 SMs x 4 schedulers x 1 instruction/cycle x 1.98e9 cycles/second
       ≈  1.05 x 10^12 instructions/second   (roughly a trillion instructions per second)

זהו throughput אדיר, ולפי חוק Little הוא בדיוק מה שמכריח מקביליות עצומה: כדי להזין קצב הנפקה כזה למרות latency של מאות מחזורים, צריך עשרות אלפי threads בטיסה בו-זמנית על פני כל ה-GPU. ברמת ה-SM הבודד: 4 schedulers שמנפיקים למרות latency של ~400 מחזורים דורשים בערך 4 x 400 = 1600 פעולות-thread בטיסה, שהן 1600/32 = 50 warps - קרוב מאוד למקסימום של 64 warps ל-SM. זו הסיבה ש-H100 תוכנן עם 64 warps תושבים ל-SM מלכתחילה: המספר נגזר ישירות מחוק Little, כדי שיהיה מספיק חומר להסתיר את ה-latency ולהזין את קצב ההנפקה.

הגשר ל-occupancy¶

נחבר את הכל לפני שנסכם. הoverhead היא הצוואר השלישי - ה-GPU בטל, ממתין לעבודה. חלק מה-latency שגורם לבטלה הזו מגיע מבחוץ (הhost, הרשת), ואת זה מתקנים עם גרפים, מיזוג וקיבוץ. אבל חלק גדול מה-latency הוא פנימי - loads מהזיכרון, שרשראות תלות - ואת זה מסתירים עם מקביליות. חוק Little הוא הכלי שמכמת בדיוק כמה מקביליות צריך: concurrency = latency x throughput, ובמונחי warps, בערך latency/32. אם יש לכם מספיק warps תושבים - ה-latency מוסתר וה-throughput נשמר. אם אין מספיק - ה-warp scheduler נשאר בלי warp כשיר להנפיק, יחידות הביצוע יושבות בטל, וה-latency הפנימי מתגלה כoverhead גלויה שמורידה אתכם מתחת ל-Roofline.

וכאן בדיוק מתחיל פרק 8. השאלה "כמה warps תושבים יש לי בפועל על ה-SM" היא בדיוק ההגדרה של occupancy, וכמה משאבים (אוגרים, shared memory) כל block צורך קובע כמה warps נכנסים. חוק Little אומר כמה warps צריך; occupancy אומר כמה warps יש. בשיעור 8.1 נחשב occupancy תיאורטי מתצורת הlaunch ומגבלות ה-H100, ונראה מדוע - כפי שכבר רמזנו - יותר occupancy הוא לא תמיד טוב יותר: צריך רק מספיק כדי לספק את חוק Little, לא יותר.

סיכום¶

תקורה (overhead) היא הזמן שבו ה-GPU יושב בטל וממתין לעבודה, בניגוד לצוואר-בקבוק (compute-bound או memory-bound) שבו המשאב המגביל עובד במלוא הכוח; זו הקטגוריה השלישית של העבודה (Compute / Memory / Overhead).
הoverhead איננה מופיעה על ה-Roofline, כי ה-Roofline מצייר רק רוחבי-פס ותפוקות ולא latency של המתנה; לכן kernel overhead-bound יושב עמוק מתחת לשני הגגות.
הhost overhead היא בערך 10 מיקרו-שניות לכל launch kernel (וה-framework מוסיף עוד מיקרו-שניות רבות של dispatch); היא הופכת דומיננטית כשמשגרים הרבה kernels קטנים, כמו כ-480 launches לטוקן ב-decode של LLM.
מקטינים host overhead על-ידי פחות launches גדולות יותר: גרפי CUDA (שיעור 5.5), מיזוג (fusion), וקיבוץ (batching).
הoverhead תקשורת (CPU-GPU על PCIe, GPU-GPU על NVLink) היא צורה של memory-boundedness על-פני מערכת מבוזרת: אותו ניתוח Roofline חל, כשרוחב-פס הרשת הוא הגג המשופע.
הסתרת latency (latency hiding) שומרת על throughput גבוה למרות latency ארוך של פעולה בודדת, על-ידי שזירת warps רבים: כש-warp נתקע, ה-scheduler מנפיק מ-warp כשיר אחר, ומיתוג ה-warp עולה בערך מחזור אחד; warps, לא threads, הם יחידת ההנפקה.
קבוצות warps שונות יכולות לתפוס pipes שונים בו-זמנית (Tensor Core, CUDA Core, LSU), כך שכל אחד מסתיר את ה-latency של האחרים - הקבלה ל-SMT בקנה-מידה גדול.
חוק Little קובע concurrency = latency x throughput: latency של 400 מחזורים ב-throughput של פקודה למחזור דורש 400 פעולות בטיסה; ב-10 פקודות למחזור - 4,000. throughput גבוה יותר דורש מקביליות גבוהה יותר באופן פרופורציוני.
ממירים latency למספר warps בחלוקה ברוחב ה-warp: שרשרת תלויה של 416 מחזורים דורשת בערך 416/32 ≈ 13 warps להסתרה מלאה (occupancy של כ-20% ב-H100); ILP הוא מנוף חלופי שמפחית את מספר ה-warps הדרוש.
תוצאת Volkov: צריך בערך 30 warps להסתיר latency של זיכרון מול בערך 24 להסתיר latency אריתמטי - דומה, למרות ש-latency הזיכרון גדול בהרבה, כי לזיכרון latency גבוה אך throughput נמוך ולאריתמטיקה להפך, והמכפלה (concurrency) יוצאת דומה.
חוק Little הוא היסוד התיאורטי ל-occupancy: הוא אומר כמה warps צריך כדי להסתיר את ה-latency; פרק 8 יראה כמה warps תושבים יש בפועל, ומדוע צריך רק מספיק ולא יותר.