לדלג לתוכן

1.4 ה Warp Scheduler ומנוע ההרצה הרצאה

בשיעור 1.3 פירקנו את ה-SM (מעבד רב-זרמי - Streaming Multiprocessor) לpipes ההרצה שלו: ה-CUDA Core (ליבת CUDA) לחשבון סקלרי, ה-SFU (יחידת פונקציות מיוחדות) לפונקציות טרנסצנדנטיות, וה-LSU (יחידת טעינה ואחסון) לגישות זיכרון. ראינו שכל אחד מהם הוא "pipe" שנתונים נכנסים אליו ותוצאות יוצאות ממנו, אבל השארנו שאלה פתוחה: מי מחליט מה נכנס לpipes בכל רגע נתון. השיעור הזה עונה על השאלה. הרכיב שמנצח על התזמורת הוא ה-warp scheduler (מתזמן ה-warp - קבוצת 32 threads), והוא הלב הפועם של ה-SM. נראה שהוא בוחר warp אחד להריץ בכל מחזור שעון, שהוא מחליף warps בעלות של אפס-כמעט-מוחלט, ושדווקא ההחלפה הזולה הזו היא הסוד מאחורי כל יכולת ה-latency hiding (הסתרת ה-latency / latency) של ה-GPU. בשיעור 1.2 הבחנו בין מקביליות (parallelism) לבין מקבילות (concurrency); כאן נראה את המנגנון החומרתי המדויק שממיר את השנייה לראשונה. זו גם התשתית שעליה ייבנו הסתרת ה-latency בפרק 7 ומצבי ה-warp בפרק 8.

מי מפעיל את הpipes - the warp scheduler

ה-warp scheduler הוא רכיב בתוך ה-SM שתפקידו, בניסוח של הגלוסרי, "להחליט איזו קבוצת threads תרוץ בכל מחזור שעון". קבוצות ה-threads האלה הן ה-warps. בכל מחזור המתזמן בוחר warp אחד מתוך מאגר ה-warps שיושבים על ה-SM, שולף את ההוראה הבאה שלו, ומנפיק (issue) אותה אל pipe ההרצה המתאים - חשבון ל-CUDA Core, טרנסצנדנטי ל-SFU, גישת זיכרון ל-LSU. זכרו את המסגור מ-1.3: הpipes עצמם הם טיפשים, הם רק מבצעים. ה-warp scheduler הוא ה"מוח" שמחליט, ולכן הוא (יחד עם קובץ האוגרים) מה שהופך את ה-SM למקבילה האמיתית של ליבת CPU, ולא הpipe הבודד.

ב-H100 יש ארבעה warp schedulers בכל SM, אחד לכל תת-מחיצה (sub-partition) של ה-SM. חלוקת המשאבים נקייה להפליא: 128 ליבות ה-FP32 של ה-SM מתחלקות לארבע, כלומר 32 ליבות FP32 לכל תת-מחיצה - בדיוק רוחב של warp אחד. לכן כשמתזמן מנפיק הוראת FP32 של warp בודד, 32 ה-threads שלו נופלים בול על 32 ליבות ה-FP32 של המחיצה שלו, ומבוצעים ב-lockstep (צעד-נעול) תחת אותה הוראה - זהו בדיוק מודל ה-SIMT (Single Instruction, Multiple Threads) שראינו.

          one SM of H100  (bird's-eye view, four sub-partitions)
  +-------------------------------------------------------------+
  |  register file: 65,536 32-bit registers (256 KiB)           |
  +------------------+------------------+------------------+-----+
  |  sub-partition 0 |  sub-partition 1 |  sub-partition 2 | 3   |
  |  warp scheduler  |  warp scheduler  |  warp scheduler  | ... |
  |  + dispatch      |  + dispatch      |  + dispatch      |     |
  |  32x FP32        |  32x FP32        |  32x FP32        |     |
  |  16x INT32/FP64  |  16x INT32/FP64  |  16x INT32/FP64  |     |
  |  1x Tensor Core  |  1x Tensor Core  |  1x Tensor Core  |     |
  |  SFU + LSU       |  SFU + LSU       |  SFU + LSU       |     |
  +------------------+------------------+------------------+-----+
  |         L1 data cache / shared memory: 256 KiB              |
  +-------------------------------------------------------------+

מכאן נובע מספר יסוד שכדאי לזכור: כל מתזמן מנפיק הוראה אחת של warp אחד למחזור, כלומר 32 threads למחזור. ארבעה מתזמנים נותנים 128 threads למחזור לכל SM. נכפיל ב-132 ה-SM של H100, ונקבל 132 x 128 = 16,896 threads שמונפקים בכל מחזור שעון ברחבי ה-GPU כולו. שימו לב שזה בדיוק מספר ליבות ה-FP32 של H100 (16,896) - לא צירוף מקרים, אלא אותו החשבון משני כיוונים.

החלפת הקשר באפס overhead - zero-overhead context switching

הטענה המרכזית של השיעור: ה-warp scheduler יכול להחליף את ה-warp הפעיל בכל מחזור שעון - בערך ננו-שנייה אחת. warp שהנפיק הוראה ונתקע ממתין לתוצאה מוחלף מיד ב-warp אחר שההוראה הבאה שלו מוכנה לרוץ. אין כאן שום "שמירת מצב" ו"טעינת מצב" כמו שאנחנו מכירים ממערכות הפעלה. ההחלפה חינם. השאלה החשובה היא למה היא חינם, ולתשובה שני חלקים.

סיבה 1 - האוגרים כבר מוקצים, אין העברת נתונים. בשיעור 1.3 ראינו שקובץ האוגרים (register file) של ה-SM ענק: 65,536 אוגרים של 32-bit לכל SM ב-H100 (סה"כ 256 KiB). כל thread שיושב על ה-SM מקבל את האוגרים הפרטיים שלו מראש, נחתך לו נתח קבוע מתוך קובץ האוגרים הגדול, וכל ה-warps התושבים חיים בו-זמנית בקובץ הזה. לכן החלפת warp לא מזיזה שום בית: המתזמן פשוט מפנה את ההוראה הבאה לקרוא מקבוצת אוגרים אחרת שכבר יושבת שם. אין spill (שפיכה) של אוגרים לזיכרון ואין טעינה מחדש - זו המהות של zero-overhead context switching.

סיבה 2 - ה-L1/shared מנוהל בתוכנה ומשותף. בניגוד ל-CPU, שבו ה-cache הוא שקוף ומנוהל בחומרה, ה-L1 cache של ה-GPU יכול להיות מנוהל-תוכנה (programmer-managed) במלואו - זהו ה-shared memory (זיכרון משותף), וראינו ב-1.3 שהוא ממופה על אותו SRAM פיזי של ה-L1. חשוב מכך, הוא משותף לכל ה-warps שמתוזמנים יחד על אותו SM. לכן החלפת warp לא גורמת ל-cache thrash (הצפת cache): ה-warp שנכנס לא "מגרש" את הנתונים של ה-warp שיצא, כי הם חולקים את אותו מרחב מנוהל. ב-CPU, לעומת זאת, כל החלפת thread נוטה למלא מחדש את ה-cache עם נתוני ה-thread החדש - וזה חלק ניכר מהמחיר.

שתי הסיבות יחד מסבירות את המספר: החלפה בערך מחזור אחד. בתדר boost של כ-1.98 GHz מחזור הוא כחצי ננו-שנייה, והגלוסרי מעגל את זה ל"בערך ננו-שנייה".

  warp switch (GPU)                     thread switch (CPU)
  -----------------                     ------------------
  the next instruction reads from       saving registers to the stack,
  another register group already        switching page tables, flushing the
  resident in the register file.        pipeline, refilling the cache.
  0 data transfers.                     hundreds to thousands of cycles.
  ~1 cycle (~1 ns)                      ~microsecond

למה זה שונה מ-CPU - fine-grained SMT

מה שה-warp scheduler עושה אינו רעיון חדש לגמרי. זהו בדיוק ה-SMT (Simultaneous Multi-Threading), הידוע בשם המסחרי "hyper-threading" ב-CPU של Intel: החזקת מספר threads "חמים" על אותה ליבה, כדי שכשאחד נתקע האחר ירוץ. ההבדל הוא בקנה המידה ובגרנולריות. ב-CPU מודרני יש שני threads חומרתיים לליבה; ב-SM של H100 יכולים לשבת עד 64 warps (2048 threads) בו-זמנית. ב-CPU ההחלפה בין threads-על מלאים (כאלה שמערכת ההפעלה מתזמנת) עולה מאות עד אלפי מחזורים, בערך מיקרו-שנייה; ב-GPU היא עולה מחזור בודד. זהו SMT דק-גרעין (fine-grained) שנלקח לקיצוניות, ויחידת התזמון היא ה-warp בן 32 הthreads, לא ה-thread הבודד.

מאפיין החלפת thread ב-CPU החלפת warp ב-GPU
מי מתזמן מערכת ההפעלה (בעיקר) ה-warp scheduler בחומרה
יחידת תזמון thread בודד warp (32 threads)
מצב שנשמר/נטען רגיסטרים, מפות דפים, ועוד שום דבר - האוגרים כבר תושבים
גורל ה-cache לרוב מתמלא מחדש (thrash) משותף, אין thrash
עלות טיפוסית מאות-אלפי מחזורים (~מיקרו-שנייה) ~מחזור אחד (~ננו-שנייה)
כמה "חמים" בו-זמנית 1-2 לליבה עד 64 warps ל-SM

מכאן, פי יותר מ-1000 מהיר יותר. וזה לא פרט טכני שולי אלא הבסיס לכל פילוסופיית ה-GPU: "throughput על פני latency". במקום להשקיע טרנזיסטורים בהקטנת ה-latency של פעולה בודדת (כמו שה-CPU עושה עם חיזוי ענפים וביצוע לא-לפי-הסדר), ה-GPU מציף את עצמו בעבודה מקבילה, ומסתיר את ה-latency מאחורי החלפות זולות.

הסתרת ה-latency - latency hiding

עכשיו נחבר הכל. הסיבה שכל המנגנון קיים היא הסתרת latency. תרחיש: warp מנפיק טעינה מ-global memory (זיכרון גלובלי) דרך ה-LSU. הטעינה יורדת אל ה-HBM3 מחוץ לשבב, וה-latency שלה עצום - מאות מחזורים. ה-warp הזה עכשיו תקוע (stalled): הוא אינו יכול להנפיק את ההוראה הבאה שלו, כי היא תלויה בערך שעדיין לא חזר. במקום להשאיר את הpipes בטלים, ה-warp scheduler עושה מיד את הדבר היחיד ההגיוני: הוא בוחר warp אחר, מוכן (eligible), שההוראה הבאה שלו כן זמינה, ומנפיק אותה. ובמחזור הבא עוד אחד. עם מספיק warps תושבים, תמיד יש warp מוכן, והpipes אף פעם לא מתבטלים.

  timeline of a single scheduler (each cell = clock cycle). W0..W3 are warps.
  L = issuing a load to global memory (hundreds of cycles until the value returns)
  . = that warp stalled, waiting for the result
  # = a ready arithmetic instruction that was issued

  time -> 1   2   3   4   5   ...        401
  W0:    [L] [.] [.] [.] [.]  .....  [#]   <- the value returned, W0 ready again
  W1:     -  [#] [L] [.] [.]  .....
  W2:     -   -  [#] [#] [L]  .....
  W3:     -   -   -  [#] [#]  .....
         ------------------------------------
  issued: W0  W1  W2  W3  W2  ...            the scheduler never idles a cycle

שימו לב לתובנה: המתזמן לא "מחכה" ל-W0. ברגע ש-W0 הנפיק את הטעינה ונתקע, המתזמן עבר הלאה. ה-latency של הטעינה של W0 "מוסתר" מאחורי העבודה המועילה של W1, W2, W3 ושאר ה-warps. זו בדיוק הסיבה ש-occupancy (תפוסה) חשובה: occupancy היא, בגסות, כמה warps תושבים על ה-SM ביחס למקסימום. תפוסה גבוהה משמעה מאגר גדול של warps שהמתזמן יכול לבחור מהם, ולכן הסתרת latency טובה יותר. את הקשר המדויק בין occupancy לביצועים, ואת מצבי ה-warp (מוכן / תקוע / פעיל) שהמתזמן עוקב אחריהם, נפרוש בפרק 8; את מכניקת הסתרת ה-latency נעמיק בשיעור 7.4. כאן די לנו בעיקרון.

מציאות ה-latency של הוראה - instruction latency

יש כאן מלכוד שמתחילים נופלים בו: הם מניחים שהתוצאה של הוראה זמינה במחזור הבא. היא לא. בניסוח של הגלוסרי, "כשמונפקת ל-warp הוראה, התוצאות בדרך כלל אינן זמינות בתוך מחזור שעון אחד, ולכן לא ניתן להנפיק הוראות תלויות". כלומר בין הנפקת הוראה לבין הרגע שתוצאתה מוכנה יש latency רב-מחזורי, ואי אפשר להנפיק הוראה שתלויה בה בזה-אחר-זה (back-to-back). שני מקורות עיקריים ל-latency הזה:

  • גישות ל-global memory - הכבד מכולם. טעינה מ-HBM שלא נמצאת ב-cache עולה מאות מחזורים (סדר גודל של 400-800), כי היא יוצאת מהשבב.
  • חלק מפעולות החשבון - גם הוראת FMA (fused multiply-add) על ליבת CUDA אינה מיידית; ה-latency שלה הוא סדר גודל של כמה מחזורים בודדים (לרוב כ-4-6). פעולות טרנסצנדנטיות על ה-SFU יקרות עוד יותר.

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

חוק Little בזעיר אנפין. נעשה אומדן. נניח latency של טעינה = 400 מחזורים, ונניח שמתזמן מנפיק הוראה אחת למחזור (throughput = 1 הוראה/מחזור). כמה הוראות עצמאיות צריכות להיות "בטיסה" בו-זמנית כדי שהמתזמן לא יתבטל אף מחזור לאורך כל חלון ה-400? חוק Little אומר: מספר-בטיסה = latency x throughput = 400 x 1 = 400 הוראות עצמאיות. אם כל warp תורם הוראה עצמאית אחת בכל רגע (כי הבאה אחריה תלויה בה), צריך בערך 400 warps כדי להסתיר את הטעינה במלואה - למתזמן בודד.

וכאן מגיע האילוץ: ל-SM של H100 יש לכל היותר 64 warps תושבים, כלומר בערך 16 warps למתזמן. 16 רחוק מ-400. המסקנה המפוכחת: הסתרת latency של זיכרון אינה מושגת מ-warps בלבד אם כל warp מנפיק רק הוראה עצמאית אחת. שני מנופים סוגרים את הפער - העלאת occupancy (עוד warps, עד התקרה), והעלאת ILP (Instruction-Level Parallelism - מקביליות ברמת ההוראה: לתת לכל warp כמה הוראות עצמאיות ברצף, למשל להנפיק מספר טעינות עצמאיות לפני שצורכים תוצאה כלשהי). לחשבון, לעומת זאת, שה-latency שלו הוא כ-4-6 מחזורים, מספיקים 4-6 warps מוכנים למתזמן - יעד קל להשגה. לכן ה-latency של הזיכרון הוא הקשה, וזה בדיוק מה שהופך kernels רבים ל-memory-bound.

הקשר לקובץ האוגרים - register pressure

נשוב לסיבה 1 להחלפה החינמית: האוגרים מוקצים מראש לכל thread תושב. זו בדיוק החרב הפיפיות. קובץ האוגרים סופי - 65,536 אוגרים ל-SM ב-H100 - וכל ה-warps התושבים חייבים להיכנס בו יחד. לכן ככל שכל thread משתמש ביותר אוגרים, כך פחות threads (ופחות warps) יכולים להיות תושבים בו-זמנית. זה נקרא register pressure (לחץ אוגרים), והוא ה-tradeoff המרכזי ששולט ב-occupancy.

החשבון פשוט: מספר ה-threads התושבים המקסימלי = min(2048, 65,536 / אוגרים-לכל-thread). הנה טבלת עבודה ל-H100 (התקרה 2048 threads = 64 warps ל-SM):

אוגרים ל-thread threads תושבים warps תושבים occupancy
32 2048 (התקרה) 64 100%
40 1638 -> 1536 (עיגול) 48 75%
64 1024 32 50%
128 512 16 25%
255 (המקסימום) 256 8 12.5%

(העיגולים נובעים מכך שהאוגרים מוקצים בגרנולריות של warp שלם ובמנות קבועות; לכן 40 אוגרים ל-thread לא נותנים בדיוק 1638 threads אלא מתעגל כלפי מטה למספר warps שלם.)

התובנה: kernel "כבד באוגרים" - כזה עם הרבה משתנים חיים בו-זמנית, לולאות פרושות (unrolled), או חישוב מורכב - מוריד את מספר ה-warps התושבים, מכווץ את מאגר הבחירה של המתזמן, ולכן פוגע בהסתרת ה-latency. יש כאן מתח אמיתי: יותר אוגרים ל-thread נותנים יותר ILP לכל warp (מנוף אחד להסתרת latency), אבל פחות warps (המנוף השני). האיזון בין השניים הוא לב אופטימיזציית ה-occupancy, ונחזור אליו בפרק 8. הcompiler nvcc מחשב את מספר האוגרים בזמן compilation, ואפשר לראות אותו עם -Xptxas -v ולהגביל אותו עם -maxrregcount או עם __launch_bounds__ - נעשה זאת בתרגול.

מקביליות מול מקבילות - the 1.2 payoff

בשיעור 1.2 הגדרנו שני מושגים שנשמעים דומים אך שונים: concurrency (מקבילות) היא ניהול של הרבה משימות שמתקדמות לסירוגין, ו-parallelism (מקביליות) היא ביצוע ממשי של הרבה משימות בו-זמנית. עכשיו יש לנו את החומרה המדויקת שממירה את הראשונה לשנייה.

ה-64 warps שיושבים על SM הם ה-concurrency: הם כולם "חיים", אבל מתזמן בודד מנפיק רק אחד מהם בכל מחזור - הוא מנהל אותם לסירוגין (interleaving), בדיוק כמו מערכת הפעלה שמנהלת הרבה תהליכים על ליבה אחת. ה-parallelism האמיתי הוא בשתי רמות: בתוך ה-warp, 32 ה-threads רצים ממש בו-זמנית על 32 הליבות (זו המקביליות של ה-SIMT); ובין תת-המחיצות, ארבעת המתזמנים של ה-SM (וכל 132 ה-SM) מנפיקים בו-זמנית. אז ה-GPU הוא גם מקבילי (מה שרץ באמת ברגע נתון - 16,896 threads ברחבי H100) וגם עמוס-מקבילות (מה שתושב וזמין - עד 2048 threads ל-SM, סה"כ מאות אלפים). ה-warp scheduler הוא הגשר: הוא לוקח את מאגר ה-concurrency הגדול ומזין ממנו את יחידות ה-parallelism הצרות בכל מחזור, כך שהן אף פעם לא רעבות. זו התמונה השלמה של "throughput על פני latency".

סיכום

  • ה-warp scheduler הוא הרכיב בתוך ה-SM שבוחר warp אחד בכל מחזור שעון ומנפיק את הוראתו לpipes ההרצה (CUDA Core, SFU, LSU) שראינו ב-1.3; הוא המוח שמנצח על הpipes הטיפשים.
  • ב-H100 יש ארבעה warp schedulers ל-SM, אחד לתת-מחיצה עם 32 ליבות FP32; כל מתזמן מנפיק 32 threads למחזור, כלומר 128 threads ל-SM למחזור, ו-16,896 threads למחזור בכל ה-GPU.
  • החלפת warp עולה בערך מחזור אחד (~ננו-שנייה) והיא "אפס-overhead" משתי סיבות: האוגרים של כל thread תושב מוקצים מראש בקובץ האוגרים ולכן ההחלפה לא מזיזה נתונים, וה-L1/shared מנוהל-תוכנה ומשותף ל-warps ולכן אין cache thrash.
  • זהו SMT דק-גרעין ("hyper-threading") שנלקח לקיצוניות בגרנולריות של warp; החלפת thread ב-CPU עולה מאות-אלפי מחזורים (~מיקרו-שנייה), כלומר ה-GPU מהיר בהחלפה פי יותר מ-1000.
  • הסתרת ה-latency היא כל התכלית: כשה-warp המנפיק נתקע על טעינה ארוכה מ-global memory, המתזמן עובר מיד ל-warp מוכן אחר; עם מספיק warps תושבים הpipes אף פעם לא מתבטלים - ולכן occupancy חשובה (פרק 8).
  • תוצאה של הוראה בדרך כלל אינה זמינה במחזור הבא, ולכן אי אפשר להנפיק הוראות תלויות זו-אחר-זו; מקורות ה-latency הם גישות global memory (מאות מחזורים) וחלק מפעולות החשבון (כמה מחזורים).
  • חוק Little בקירוב: להסתרת latency של 400 מחזורים בקצב הנפקה של הוראה אחת למחזור צריך כ-400 הוראות עצמאיות בטיסה; מכיוון שמתזמן מחזיק לכל היותר כ-16 warps, מסתירים latency של זיכרון בשילוב של occupancy גבוהה ו-ILP (כמה הוראות עצמאיות ל-warp).
  • ההחלפה חינם רק כי האוגרים תושבים מראש, ומכאן ה-tradeoff: שימוש כבד באוגרים לכל thread (register pressure) מקטין את מספר ה-warps התושבים ופוגע בהסתרת ה-latency; ב-H100, 32 אוגרים/thread נותנים 100% occupancy, ואילו 128 אוגרים/thread מורידים ל-25%.
  • ה-warp scheduler הוא הגשר מ-1.2 בין concurrency (עד 2048 threads תושבים ל-SM, מנוהלים לסירוגין) ל-parallelism (16,896 threads שרצים ממש בכל מחזור): הוא ממיר מאגר גדול של עבודה זמינה לתפוקה רציפה בpipes הצרים.