1.2 ה SM לעומק פתרון
פתרון - ה-SM לעומק¶
כאן נעבור על הפתרון המלא של כל תרגיל. שימו לב: הפלטים המספריים בהמשך הם דוגמאות מהרצה על H100 SXM, והם ישתנו בין כרטיסים - זו כל מטרת התרגילים המעשיים. המבנה, ההיגיון וההודעות זהים בכל מכונה. אם אתם על T4 ב-Colab, החליפו את -arch=sm_90a ב--arch=sm_75, ואת גרסת ה-compute_70 בתרגיל 5 בגרסה ישנה מזו של הכרטיס שלכם.
פתרון תרגיל 1 - קריאת המאפיינים של ה-SM שלכם¶
הקומפילציה וההרצה:
הפלט על H100 SXM:
GPU : NVIDIA H100 80GB HBM3
Compute Capability : 9.0
SMs : 132
Warp size : 32
Max threads / SM : 2048
Max warps / SM : 64
Registers / SM : 65536
Shared mem / SM : 228 KiB
ההצלבה עם nvidia-smi:
ה-compute capability הוא 9.0, ולפי טבלת המשפחות בהרצאה גרסה ראשית 9 שייכת לHopper. שני המקורות (ה-API וה-nvidia-smi) מדווחים 9.0 - התאמה.
למה זה עבד: cudaGetDeviceProperties ממלא את המבנה cudaDeviceProp ישירות מהחומרה בזמן ריצה, ולכן הוא תמיד מדויק לכרטיס שעליו רצים - עדיף על שינון מספרים. שימו לב לפער קטן שראוי להסביר: השדה sharedMemPerMultiprocessor מדווח 228 KiB, בעוד שבהרצאה אמרנו 256 KiB. אין כאן סתירה: 256 KiB הוא ה-L1/shared המאוחד הפיזי של ה-SM, וה-API מדווח את הפלח שניתן להקצות כ-shared memory (רוב ה-256, אך לא הכל - חלק שמור ל-L1). מספר ה-warps חושב כ-maxThreadsPerMultiProcessor / warpSize = 2048 / 32 = 64.
איך להכליל: התוכנית הזו היא כלי הבסיס לחקר כל כרטיס חדש שנתקלים בו. אותם שדות (multiProcessorCount, maxThreadsPerMultiProcessor, regsPerMultiprocessor, warpSize) הם הקלט לכל חישוב occupancy בהמשך הקורס. על כרטיס אחר תקבלו מספרים אחרים אך אותו מבנה בדיוק - למשל T4 יראה 7.5, 40 SMs ו-64 warps ל-SM.
פתרון תרגיל 2 - ה-SM מול ה"ליבה": מדוע ה-SM הוא המעבד¶
שלב 1 - pipe מול מעבד. הpipe הוא יחידה פונקציונלית שיודעת רק פעולה אחת: לקבל אופרנדים ולהחזיר תוצאה מומרת. חסרים לליבת CUDA בודדת שלושה דברים שהופכים משהו למעבד: מונה-פקודות משלה (היא לא בוחרת מה להריץ), מתזמן משלה (מישהו אחר מנפיק לה את הפקודה), וזיכרון פרטי מעבר לאוגרים שמוזנים אליה. היא לא מחליטה ולא זוכרת - היא רק ממירה.
שלב 2 - שלושת המרכיבים של מעבד שקיימים ב-SM:
| מרכיב | מה הוא תורם |
|---|---|
| קובץ אוגרים (register file) | מחזיק את מצב כל הthreads החיים - הזיכרון של המעבד |
| ליבות (cores) | ממירות את הנתונים - יחידת החישוב בפועל |
| מתזמני warp (warp schedulers) | מחליטים בכל מחזור איזו קבוצת threads תרוץ ומנפיקים לה פקודה - הבקרה |
שלושתם יחד = מעבד עצמאי. ליבת CUDA לבדה מחזיקה רק את השני מהם.
שלב 3 - התשובה לחבר. "16,896 CUDA Cores" הוא מספר נתיבים (lanes), לא מספר מעבדים. ליבת CUDA היא pipe בלי מתזמן ובלי מונה-פקודות משלה, ולכן היא אינה יכולה להריץ תוכנית עצמאית - היא מבצעת רק את הפקודה שמתזמן ה-warp מנפיק לקבוצה שלמה של 32 נתיבים בבת אחת (מודל SIMT). המספר שמתאר "כמה יחידות עצמאיות" יש בכרטיס הוא מספר ה-SMs - 132 ב-H100. כל SM הוא מעבד עצמאי; כל "core" בתוכו הוא רק נתיב.
למה זה עבד: ההבחנה כולה נשענת על השאלה "מי מחליט מה לרוץ". מעבד מחליט בעצמו (יש לו מתזמן ומונה-פקודות); pipe לא מחליט (מוזן מבחוץ). ה-SM מחליט - הוא היחידה שמתזמנת. הליבה רק מבצעת.
איך להכליל: בכל פעם שתראו מספר שיווקי גדול של "cores", תרגמו אותו מיד ל"נתיבים" ותשאלו כמה SMs יש. זה מתקן את המודל המנטלי מ"אלפי מעבדים קטנים" ל"מאה ומשהו מעבדים, כל אחד רחב מאוד". אותה הבחנה תחזור כשנשווה CUDA Cores ל-Tensor Cores - שניהם pipes, נבדלים בסוג ההמרה.
פתרון תרגיל 3 - חישוב parallel מול concurrent לכרטיס שלכם¶
עבור H100 (132 SMs, 2048 threads תושבים ל-SM):
parallel (truly parallel):
4 sub-partitions x 32 threads x 132 SM = 16,896 threads per cycle
concurrent (resident):
2048 threads/SM x 132 SM = 270,336 threads
ratio:
270,336 / 16,896 = 16 (also, per SM: 2048 / 128 = 16)
שלב 4 - תפקיד הפער. הפער בין 128 הthreads שמתקדמים בכל מחזור לבין 2,048 התושבים הוא מאגר הסתרת ה-latency. כשה-warps הרצים כרגע נתקעים בהמתנה ל-global memory (latency של מאות מחזורים), מתזמן ה-warp שולף מהמאגר warp מוכן אחר וממשיך לעבוד - הpipes לא יושבים בטל. אילו ה-SM היה יכול להחזיק בדיוק 128 threads (יחס 1:1), לא היה מאגר: ברגע שאותם 128 threads היו ניגשים לזיכרון, ה-SM היה משתתק לחלוטין למאות מחזורים, כי אין למי להחליף. הפער הוא בדיוק מה שקונה את הסבילות ל-latency.
למה זה עבד: ה"parallel" bound ברוחב ההנפקה של החומרה (מתזמן אחד x warp אחד x 4 תת-מחיצות = 128 ל-SM), ואילו ה"concurrent" bound בכמה מצב אפשר להחזיק בקובץ האוגרים (2,048 threads). שני חסמים שונים לגמרי, ומכאן שני מספרים בסדרי גודל שונים.
איך להכליל: היחס משתנה מעט בין ארכיטקטורות, אבל התבנית קבועה: תמיד יש הרבה יותר threads תושבים ממתקדמים, והפער הזה הוא הדלק של הסתרת ה-latency. כשנלמד occupancy נראה שהמטרה המעשית היא למלא את המאגר הזה - להעמיס מספיק warps תושבים כדי שתמיד יהיה למתזמן למי להחליף.
פתרון תרגיל 4 - מדוע "אין חיזוי הסתעפויות" מקובל ב-GPU אבל לא ב-CPU¶
שלב 1 - מה עושה חיזוי הסתעפויות ב-CPU. כשה-CPU מגיע ל-if, התנאי לעיתים עדיין מחושב (למשל מחכה לערך מהזיכרון). כדי לא לעצור את הpipeline העמוק, ה-CPU מנחש לאיזה ענף ילך ומתחיל להריץ אותו ספקולטיבית. אם ניחש נכון - הרוויח. אם טעה - צריך לרוקן את הpipeline (pipeline flush): לזרוק את כל העבודה הספקולטיבית ולהתחיל מחדש מהענף הנכון, מחיר של עשרות מחזורים.
שלב 2 - למה ל-CPU אין ברירה. ל-CPU יש מעט threads ליבה. כשהthread היחיד שרץ על הליבה תקוע בהמתנה לתנאי, אין לליבה thread אחר מוכן להחליף אליו - אם לא תנחש, הpipeline פשוט יעמוד ריק. לכן הניחוש, גם עם סיכון לריקון יקר, עדיף על לעמוד בטל. ה-CPU קונה ביצועים של thread בודד במחיר סיליקון וסיכון.
שלב 3 - איך ה-GPU פותר בלי לנחש. למתזמן ה-warp ב-SM יש עשרות warps תושבים ומוכנים. כשה-warp הנוכחי מגיע להסתעפות שהתנאי שלה עדיין לא מוכן, המתזמן פשוט מחליף ל-warp אחר שהאופרנדים שלו מוכנים בעלות של מחזור אחד, וחוזר ל-warp הראשון כשהתנאי יהיה מוכן. אין ניחוש, אין ספקולציה, ואין ריקון pipeline.
שלב 4 - המשפט המסכם. ל-GPU יש בשפע משאב שאין ל-CPU: מקביליות בצורת warps תושבים מוכנים. בגלל שתמיד יש למי להחליף, ה-GPU יכול להסתיר את זמן ההמתנה על ידי החלפה במקום על ידי ניחוש. ל-CPU, עם מעט threads, אין את המאגר הזה, ולכן הניחוש (חיזוי הסתעפויות) הוא הכרחי. הוויתור על הניבוי מקובל ב-GPU בדיוק כי המקביליות מחליפה אותו.
למה זה עבד: גם ה-CPU וגם ה-GPU מתמודדים עם אותה בעיה - הpipeline לא יכול לעמוד בטל בזמן שתנאי מחושב. הם פותרים אותה בשני משאבים שונים: ה-CPU בספקולציה (סיליקון וניחוש), ה-GPU במקביליות (החלפת warp). כל אחד בוחר את מה שיש לו בשפע.
איך להכליל: אותו היגיון חל על כל מנגנוני ה"thread הבודד" שה-GPU ויתר עליהם - ביצוע מחוץ לסדר, ביצוע ספקולטיבי, caches ענקיים אוטומטיים. בכולם התשובה זהה: היכן שה-CPU מוציא סיליקון כדי להאיץ thread בודד, ה-GPU מוציא את אותו סיליקון על עוד יחידות חישוב ועוד warps, ומסתיר את הlatency בהחלפה. זה עובד רק כשיש מספיק מקביליות להאכיל את המאגר.
פתרון תרגיל 5 - תאימות SASS בין גרסאות ראשיות¶
שלב 2-3 - רק SASS של Volta. הקומפילציה מטמיעה רק קוד ילידי לגרסה ראשית 7:
ההרצה על H100 (גרסה ראשית 9) נכשלת, ומקרו ה-CUDA_CHECK תופס:
CUDA error at vecadd.cu:52
'cudaGetLastError()'
-> no kernel image is available for execution on the device
ה-driver מחפש בקובץ ההרצה SASS שתואם ל-Hopper, ולא מוצא - יש רק SASS של Volta, שאינו תואם בין גרסאות ראשיות, ואין PTX לגיבוי.
שלב 4-5 - PTX של Volta. עכשיו מטמיעים PTX נייד במקום SASS ילידי:
ההרצה על H100 מצליחה:
ההבדל: compute_70 בשדה ה-code מטמיע PTX (ייצוג ביניים), ולא SASS. כשה-driver לא מוצא SASS מתאים ל-Hopper, הוא מבצע ל-PTX קומפילציה בזמן ריצה (JIT) ל-SASS של Hopper, ולכן ה-kernel רץ.
למה זה עבד: SASS הוא קוד המכונה הילידי, כבול לגרסת ה-SM שאליה קומפל, ותאימות בין גרסאות ראשיות אינה מובטחת - ולכן SASS של Volta (7) פשוט אינו "תמונת kernel" חוקית עבור Hopper (9). PTX, לעומתו, הוא ייצוג ביניים נייד קדימה: ה-driver תמיד יכול לתרגם אותו לגרסת ה-SM הנוכחית. ההבחנה בין code=sm_XX (מטמיע SASS) לבין code=compute_XX (מטמיע PTX) היא בדיוק ההבדל בין תמונה כבולה-ארכיטקטורה לבין תמונה ניידת.
איך להכליל: בפרודקשן מטמיעים בדרך כלל את שניהם - -gencode arch=compute_90,code=sm_90 לביצועים מיטביים על Hopper, ועוד arch=compute_90,code=compute_90 כרשת ביטחון של PTX לכרטיסים חדשים יותר (תאימות קדימה דרך JIT). הדגל הנוח -arch=sm_90a עושה בעצם את שני הדברים. את ההבחנה המלאה בין PTX ל-SASS, ולמה ה-JIT מוסיף overhead חד-פעמית בהרצה הראשונה, נפתח בפרק ה-PTX/SASS.
פתרון תרגיל 6 (בונוס) - לחץ אוגרים וגבול ה-warps התושבים¶
שלב 1 - הנוסחה. מספר הthreads המקסימלי מבחינת אוגרים בלבד הוא רצפה(65536 / R), ומספר ה-warps הוא זה חלקי 32 (bound מלמעלה בתקרת 64 ה-warps של החומרה).
שלב 2-3 - הצבה:
R = 32 : 65536/32 = 2048 threads = 64 warps -> full (bound by the 64 cap)
R = 64 : 65536/64 = 1024 threads = 32 warps -> half
R = 128 : 65536/128 = 512 threads = 16 warps -> quarter
R = 255 : 65536/255 = 257 threads ~ 8 warps -> minimal
עד R = 32 מגיעים ל-64 ה-warps המלאים (התפוסה המקסימלית). מעל R = 32 מתחילים לאבד warps תושבים, כי קובץ האוגרים מתרוקן לפני שממלאים את תקרת ה-64. זו התמורה: יותר אוגרים לthread = פחות threads תושבים = מאגר הסתרה קטן יותר = פחות יכולת להסתיר latency.
שלב 4 - מדידת R בפועל. מקמפלים עם דיווח מפורט מ-ptxas:
ptxas info : Compiling entry function '_Z6vecAddPKfS0_Pfi' for 'sm_90a'
ptxas info : Used 16 registers, 376 bytes cmem[0]
ה-kernel משתמש ב-16 אוגרים. הצבה בנוסחה: 65536/16 = 4096 threads, אבל זה מעל תקרת ה-2048 של החומרה, ולכן ה-vecAdd מגיע ל-64 warps תושבים מלאים - האוגרים אינם צוואר הבקבוק שלו כלל.
למה זה עבד: קובץ האוגרים הוא משאב סופי שמתחלק בין כל הthreads התושבים. מכיוון שכל thread תושב מחזיק את האוגרים שלו פיזית וקבוע (זה מה שמאפשר את החלפת ההקשר באפס overhead), ככל שthread צורך יותר אוגרים, פחות threads נכנסים לאותו קובץ בגודל קבוע. vecAdd צנוע (16 אוגרים) ולכן רץ בתפוסה מלאה; kernel כבד-אוגרים היה מוגבל.
איך להכליל: זהו הקשר הישיר בין קובץ האוגרים ל-occupancy: לחץ אוגרים (register pressure) הוא אחד משלושת החסמים על התפוסה (יחד עם shared memory ומגבלת ה-blocks). כשקרנל סובל מ-occupancy נמוך, -Xptxas -v הוא הצעד הראשון לאבחון - הוא מראה אם האוגרים הם האשמים. אפשר גם להגביל אותם בכפייה עם nvcc --maxrregcount=N או __launch_bounds__, במחיר של אולי spill לזיכרון. את המסגרת המלאה - מחשבון ה-occupancy והאיזון בין המשאבים - נפתח בפרק הביצועים.