1.4 ה Warp Scheduler ומנוע ההרצה פתרון
פתרון - ה-Warp Scheduler ומנוע ההרצה¶
נעבור תרגיל אחרי תרגיל. כל המספרים והפלטים כאן הם מ-H100 (SXM) עם CUDA 12.4 ותדר boost של 1.98 GHz; מספרי האוגרים המדויקים שה-ptxas מדווח וזמני ה-ncu אצלכם עשויים להשתנות מעט בין גרסאות toolkit, וזה תקין. מי שעובד על T4 יראה מספרים אחרים (64 ליבות FP32 ל-SM, תקרה של 32 warps ל-SM), ונציין זאת היכן שרלוונטי.
פתרון תרגיל 1 - קצב ההנפקה של ה-SM¶
הנתונים מ-cudaGetDeviceProperties על H100: 132 SM, warpSize 32, boost כ-1.98 GHz.
חישוב שלב-אחר-שלב:
warp scheduler מנפיק warp אחד למחזור = 32 threads/מחזור
4 מתזמנים ל-SM -> 4 x 32 = 128 threads/SM/מחזור
132 SM -> 132 x 128 = 16,896 threads/מחזור (כל ה-GPU)
מספר ליבות FP32 ל-GPU = 132 SM x 128 FP32/SM = 16,896 ליבות FP32
שני המספרים בסעיף 3 זהים: 16,896. זו אותה כמות משני כיוונים - מצד ההנפקה (4 מתזמנים x 32 threads) ומצד ההרצה (128 ליבות FP32). זה לא מקרי: 128 ליבות ה-FP32 של ה-SM מתחלקות בדיוק בין ארבע תת-המחיצות, 32 לכל אחת, ולכן הנפקת warp אחד למחזור בכל מתזמן מזינה בדיוק את כל ליבות ה-FP32 שלו.
קצב ההנפקה בהוראות-thread לשנייה:
16,896 threads/מחזור x 1.98e9 מחזורים/שנייה = 3.345e13 = 33.45 T threads-instr/s
שיא FP32: 16,896 x 2 FLOP/FMA x 1.98e9 = 66.9 TFLOPS
היחס בין 66.9 TFLOPS ל-33.45 הוא בדיוק 2 - כי כל הוראת FMA היא שתי פעולות נקודה-צפה (כפל ועוד). סדר הגודל מתיישב.
למה זה עבד: קצב ההנפקה נגזר משלושה מספרים בלבד - מתזמנים ל-SM, רוחב warp, ומספר SM - כי בכל מחזור כל מתזמן מנפיק warp שלם. איך להכליל: על T4 (Turing) יש ארבעה מתזמנים ל-SM, כל תת-מחיצה עם 16 נתיבי FP32 (סה"כ 64 FP32 ל-SM), 40 SM ותדר 1.59 GHz. באותה שיטה, קצב ההנפקה הוא 40 x 4 x 32 = 5,120 threads/מחזור, בעוד ש-40 x 64 = 2,560 הוא מספר נתיבי ה-FP32 בכל ה-GPU (ולכן הנפקת warp של FP32 נמשכת שני מחזורים). תמיד גִזרו את קצב ההנפקה מספירת המתזמנים והליבות של הדור הספציפי, לא ממספר מקובע.
פתרון תרגיל 2 - למה ההחלפה זולה פי אלף¶
זהו תרגיל ניתוח. התשובה המלאה:
החלפת thread ב-CPU. מערכת ההפעלה שומרת את כל הרגיסטרים הארכיטקטוניים של ה-thread היוצא ל-stack, מחליפה את הקשר הזיכרון (page tables / הצבעת ה-TLB), מרוקנת את ה-pipeline, וטוענת את מצב ה-thread הנכנס. חמור מכך, ה-thread הנכנס מוצא cache מלא בנתונים של ה-thread הקודם, וממלא אותו מחדש תוך כדי ריצה (cache thrash). העלות הכוללת: מאות עד אלפי מחזורים, סדר גודל של מיקרו-שנייה.
החלפת warp ב-GPU. המתזמן פשוט בוחר warp אחר מהמאגר ומנפיק את הוראתו. שתי סיבות שההחלפה אינה מזיזה נתונים:
1. האוגרים של כל thread תושב מוקצים לו מראש מתוך קובץ האוגרים של ה-SM, וכל ה-warps התושבים חיים בו בו-זמנית. ההחלפה רק מפנה את ההוראה הבאה לקרוא מקבוצת אוגרים אחרת שכבר יושבת שם - אין spill ואין reload.
2. ה-L1/shared מנוהל-תוכנה ומשותף לכל ה-warps המתוזמנים יחד; ה-warp שנכנס לא מגרש את נתוני ה-warp שיצא, ולכן אין cache thrash.
היחס והמסקנה:
וזה מתחבר לקנה המידה: ב-CPU יש 1-2 threads חמים לליבה, ב-SM עד 64 warps חמים. תשובה לסעיף 4: אם היינו שומרים ומשחזרים אוגרים בכל החלפה, ההחלפה הייתה הופכת מ"שינוי מצביע" (מחזור אחד) ל"העברת עשרות אוגרים אל/מן הזיכרון" (עשרות עד מאות מחזורים), והתכונה הקריטית של החלפה בכל מחזור שעון הייתה נעלמת - וללא החלפה-כל-מחזור, אין הסתרת latency יעילה.
למה זה עבד: הזולוּת נובעת ישירות מהחלטת תכן - להקצות מראש את כל האוגרים של כל ה-warps התושבים. איך להכליל: כל מנגנון "אפס-overhead" נשען על אותו עיקרון - להימנע מהעברת מצב בזמן ההחלפה על ידי כך שהמצב כבר יושב במקומו הסופי.
פתרון תרגיל 3 - חוק Little בזעיר אנפין¶
חישוב חוק Little עבור זיכרון:
latency = 400 מחזורים, throughput = 1 הוראה/מחזור
מספר-בטיסה = latency x throughput = 400 x 1 = 400 הוראות עצמאיות
אם כל warp תורם הוראה עצמאית אחת בכל רגע, צריך כ-400 warps למתזמן. אבל מתזמן בודד ב-H100 מחזיק לכל היותר:
16 רחוק מ-400. המסקנה: עם warps בלבד, וכשכל warp מנפיק רק הוראה עצמאית אחת, אי אפשר להסתיר במלואו latency של זיכרון. שני המנופים שסוגרים את הפער (סעיף 3):
- העלאת occupancy - כמה שיותר warps תושבים, עד התקרה של 64 ל-SM.
- העלאת ILP - לתת לכל warp כמה הוראות עצמאיות ברצף (למשל להנפיק מספר טעינות עצמאיות לפני שצורכים תוצאה כלשהי), כך שמספר-בטיסה אפקטיבי = warps x הוראות-עצמאיות-לכל-warp. אם כל warp מחזיק ~25 טעינות עצמאיות בטיסה, 16 warps x 25 = 400.
חישוב עבור חשבון (סעיף 4):
כלומר כ-5 warps מוכנים למתזמן מספיקים כדי להסתיר latency של חשבון - יעד קל להשגה. לכן latency של חשבון קל להסתרה (עשרות מחזורים בסך הכל דורשים בודדי warps), ואילו latency של זיכרון (מאות מחזורים) קשה, ודורש occupancy גבוהה ו-ILP משמעותי. זו הסיבה ש-kernels רבים memory-bound.
למה זה עבד: חוק Little קושר בצורה ישירה latency, throughput וכמות המקביליות הדרושה כדי לרווי pipe. איך להכליל: אותו חישוב מנחה כל אופטימיזציית occupancy - אמדו את ה-latency הדומיננטית של ה-kernel, חלקו בקצב ההנפקה, וזה יעד המקביליות (warps x ILP) שאתם צריכים; נחזור לזה בשיעור 7.4.
פתרון תרגיל 4 - מדידת צריכת האוגרים¶
$ nvcc -arch=sm_90 -Xptxas -v regpressure.cu -o regpressure
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z5heavyPKfPfi' for 'sm_90'
ptxas info : Function properties for _Z5heavyPKfPfi
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 40 registers, 372 bytes cmem[0]
הפונקציה heavy צורכת 40 אוגרים. החישוב הידני ל-H100 (65,536 אוגרים ל-SM, תקרה 2048 threads):
threads תושבים = min(2048, 65536 / 40) = min(2048, 1638) = 1638
מעוגל כלפי מטה למספר warps שלם: 1638 / 32 = 51.2 -> 51 warps -> 1632 threads
(בפועל הקצאה בגרנולריות של מנות אוגרים ל-warp נותנת ~48 warps)
occupancy ~ 48 / 64 = 75%
הרצות עם הגבלת אוגרים:
$ nvcc -arch=sm_90 -Xptxas -v -maxrregcount=32 regpressure.cu -o regpressure
ptxas info : Function properties for _Z5heavyPKfPfi
24 bytes stack frame, 28 bytes spill stores, 28 bytes spill loads
ptxas info : Used 32 registers, 372 bytes cmem[0]
-> 65536/32 = 2048 threads = 64 warps = 100% occupancy (אבל יש spills!)
$ nvcc -arch=sm_90 -Xptxas -v -maxrregcount=64 regpressure.cu -o regpressure
ptxas info : Used 40 registers, 372 bytes cmem[0]
-> ה-kernel צריך רק 40, אז המגבלה 64 לא כובלת; נשאר 40 -> ~48 warps -> 75%
$ nvcc -arch=sm_90 -Xptxas -v -maxrregcount=128 regpressure.cu -o regpressure
ptxas info : Used 40 registers, 372 bytes cmem[0]
-> שוב לא כובל; 40 אוגרים, ~48 warps, 75%
התובנה החשובה: הגבלה ל-32 העלתה את ה-occupancy ל-100% אבל יצרה spill stores/spill loads - הcompiler שפך משתנים ל-local memory (שיושב ב-global memory האיטי). זה ממחיש שהעלאת occupancy בכוח אינה חינם: היא יכולה להחליף בעיה אחת (מעט warps) באחרת (spills איטיים). המגבלות 64 ו-128 לא השפיעו כי ה-kernel ממילא צריך רק 40.
למה זה עבד: -Xptxas -v חושף את מספר האוגרים ש-ptxas הקצה, וזה בדיוק המספר שקובע כמה warps נכנסים על SM. איך להכליל: כדי לשלוט ב-occupancy אפשר לכוון אוגרים עם -maxrregcount או, מדויק יותר, עם __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) ישירות על ה-kernel - שגורם לcompiler לכוון את מספר האוגרים ליעד ה-occupancy שביקשתם, ובלי להשפיע על kernels אחרים בקובץ.
פתרון תרגיל 5 - אימות התפוסה עם ה-API¶
int blocks = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocks, heavy, /*blockSize=*/256, /*dynamicSMem=*/0));
int warpsPerBlock = 256 / 32; // = 8
int activeWarps = blocks * warpsPerBlock;
printf("active blocks/SM = %d, active warps/SM = %d, occupancy = %.1f%%\n",
blocks, activeWarps, 100.0 * activeWarps / 64.0);
$ nvcc -arch=sm_90 regpressure.cu -o regpressure && ./regpressure
active blocks/SM = 6, active warps/SM = 48, occupancy = 75.0%
6 blocks x 8 warps = 48 warps תושבים = 75% - בדיוק כמו החישוב הידני מתרגיל 4 (40 אוגרים -> ~48 warps). ה-API וה-חישוב הידני מסכימים.
שינוי גודל ה-block:
blockSize = 128 -> 12 blocks/SM x 4 warps = 48 warps = 75%
blockSize = 512 -> 3 blocks/SM x 16 warps = 48 warps = 75%
ה-occupancy נשאר 75% כי הוא נחסם על ידי האוגרים (40 x threads מגיע ל-65,536 בסביבות 1536-1638 threads), לא על ידי גודל ה-block - כל עוד גודל ה-block מתחלק יפה. אם היינו בוחרים block שלא מתחלק בתקרות (למשל 768), היינו מפסידים warps לגרנולריות.
למה זה עבד: cudaOccupancyMaxActiveBlocksPerMultiprocessor מריץ את אותו מחשבון תפוסה של ה-driver, ששוקלל את כל האילוצים יחד (אוגרים, shared memory, warp slots) ומחזיר את הצוואר-בקבוק. איך להכליל: השתמשו בפונקציה הזו (או ב-cudaOccupancyMaxPotentialBlockSize שבוחרת גודל block אופטימלי) במקום לחשב ידנית - כך הקוד מסתגל לכל GPU בזמן ריצה ואינו תלוי במספרים קשיחים.
פתרון תרגיל 6 - מקבילות מול מקביליות, חזרה ל-1.2¶
החישובים עבור ה-launch (grid 4096, block 256):
1. concurrency פוטנציאלי (כל העבודה): 4096 x 256 = 1,048,576 threads
2. concurrency פעיל (תושבים בו-זמנית): 48 warps x 32 = 1536 threads/SM
x 132 SM = 202,752 threads
3. parallelism אמיתי (רצים כל מחזור): 128 threads/SM x 132 = 16,896 threads
הסדר מהגדול לקטן:
1,048,576 (עבודה כוללת, concurrency פוטנציאלי)
> 202,752 (תושבים בו-זמנית, concurrency פעיל)
> 16,896 (רצים ממש בכל מחזור, parallelism אמיתי)
ההסבר: המיליון threads הם ה-concurrency - עבודה שמנוהלת לסירוגין, לא רצה כולה בו-זמנית. ה-202,752 התושבים הם ה-concurrency ה"חם" שהמתזמנים בוחרים ממנו. ה-16,896 הרצים בכל מחזור הם ה-parallelism האמיתי - מה שקורה ממש ברגע נתון. ה-warp scheduler הוא הגשר: הוא לוקח את המאגר הגדול של warps תושבים וממיר אותו, warp אחד למחזור לכל מתזמן, לתפוקה רציפה על 128 יחידות ההרצה של כל SM - כך שהpipes הצרים אף פעם לא רעבים למרות שהעבודה הכוללת עצומה.
למה זה עבד: שלושת המספרים הם שלוש רמות שונות של אותו launch - נדרש, תושב, ורץ - וההפרש ביניהם הוא בדיוק המרווח שהמתזמן ממלא. איך להכליל: זו התמונה מ-1.2 בממדים אמיתיים; היחס בין העבודה הכוללת לתושבים קובע כמה "גלים" (waves) של blocks יידרשו, והיחס בין תושבים לרצים הוא ה-headroom להסתרת latency.
פתרון תרגיל 7 (בונוס) - פרופיילינג של הסתרת latency¶
$ ncu --set basic ./regpressure
heavy(const float *, float *, int), 2024-... Section: Occupancy
Achieved Occupancy % 74.8
Theoretical Occupancy % 75.0
Section: Scheduler Statistics
Issued Warp Per Scheduler 0.71
Issue Efficiency (Issued IPC) % ~71
Section: Warp State Statistics
Warp Cycles Per Issued Instruction ...
Stall (dominant) Stall Wait / Short Scoreboard
ה-Achieved Occupancy (74.8%) קרוב מאוד ל-Theoretical (75%) שחישבנו - ה-kernel אכן ממלא את מאגר ה-warps כמעט במלואו. סיבת ה-stall הדומיננטית ב-heavy היא תלות חשבון (Short Scoreboard / Wait על שרשרת ה-FMA), לא זיכרון - כי ה-kernel כמעט ולא ניגש ל-global memory אלא רק מבצע שרשרת FMA תלויה. זה תואם את תרגיל 3: latency של חשבון מוסתר בקלות עם מספיק warps.
השוואת -maxrregcount=32 (100% תאורטי, עם spills) מול ברירת המחדל (75%, בלי spills):
default (40 reg) maxrregcount=32
Theoretical Occupancy 75% 100%
Achieved Occupancy ~74.8% ~97%
spill stores/loads 0 28 + 28 bytes
זמן ריצה של ה-kernel קצר יותר לרוב איטי יותר או שווה
התובנה: העלאת ה-occupancy ל-100% על ידי כפיית 32 אוגרים הכניסה spills, וה-spills (גישות ל-local/global memory) לרוב מבטלים את הרווח - לפעמים ה-kernel אפילו איטי יותר למרות ה-occupancy הגבוה יותר. occupancy גבוהה היא אמצעי, לא מטרה.
למה זה עבד: ncu קורא מוני חומרה אמיתיים על ה-SM - warps פעילים, מחזורי הנפקה, וסיבות stall - ולכן הוא מראה את מה שבאמת קרה, לא את מה שקיווינו. איך להכליל: תמיד מדדו לפני שמכווננים occupancy; בדקו את סיבת ה-stall (זיכרון מול תלות מול הנפקה) כדי לדעת איזה מנוף למשוך, ואל תניחו שיותר occupancy = יותר מהירות. את מתודולוגיית ה-profiling המלאה נפרוש בפרק 8.