1.1 מבנה ה GPU GPC, TPC ו SM תרגול

תרגול - מבנה ה-GPU - GPC, TPC ו-SM

בתרגול הזה תחקרו את היררכיית החומרה של GPU אמיתי, במקום רק לקרוא עליה. תשאלו את הכרטיס כמה SM-ים יש לו בשלוש דרכים שונות (nvidia-smi, deviceQuery, ותוכנית CUDA קטנה משלכם), תצליבו את המספר מול ה-whitepaper כדי לשחזר את מבנה ה-GPC/TPC, תשתמשו בשרשרת ההכלה כדי להסביר למה שני block-ים בקלאסטר יכולים לחלוק זיכרון ושני block-ים שרירותיים לא, ולבסוף תשוו את ספירת ה-SM על פני שלושה דורות. עבדו לפי הסדר - כל תרגיל נשען על הקודם. הריצו על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס שלנו, אבל כל תרגיל כאן עובד על כל GPU של NVIDIA - פשוט תקבלו מספרים אחרים, וזו כל הנקודה.

הכנה

ודאו ש-nvidia-smi ו-nvcc זמינים:

nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv
nvcc --version

בררו את ה-compute capability של הכרטיס כדי לדעת באיזו ארכיטקטורה להדר (sm_90a ל-Hopper, sm_80 ל-A100, sm_75 ל-T4):

nvidia-smi --query-gpu=compute_cap --format=csv,noheader

תרגיל 1 - שלוש דרכים לספור SM-ים

1. הריצו nvidia-smi וזהו את שם הכרטיס ואת כמות הזיכרון. שימו לב: nvidia-smi הרגיל אינו מדפיס את מספר ה-SM-ים - זו כבר הפתעה ראשונה.
2. בנו וכתבו תוכנית CUDA קטנה smcount.cu שמדפיסה את שם הכרטיס, את ה-compute capability, ואת multiProcessorCount (השתמשו ב-cudaGetDeviceProperties ובמקרו CUDA_CHECK מההרצאה). הדרו עם nvcc -O2 -arch=sm_90a -o smcount smcount.cu והריצו.
3. אם deviceQuery מותקן (מ-CUDA Samples), הריצו גם אותו וחפשו את השורה (N) Multiprocessors.
4. ודאו ששלושת המקורות שראיתם עקביים, ורשמו: מה כן חשף nvidia-smi ומה לא.

רמז: nvidia-smi הוא כלי ניטור/ניהול (טמפרטורה, זיכרון, תהליכים) ולא כלי ארכיטקטורה - ולכן אין בו שדה SM. השדה multiProcessorCount ב-cudaDeviceProp הוא המקור התוכנתי הישיר למספר ה-SM-ים. אם הקומפילציה נכשלת עם שגיאת ארכיטקטורה, החליפו את sm_90a בערך שקיבלתם מ-nvidia-smi --query-gpu=compute_cap.

תרגיל 2 - שחזור ה-TPC וההצלבה מול ה-whitepaper

1. קחו את מספר ה-SM-ים שקיבלתם בתרגיל 1 (על H100: 132). מכיוון ש-TPC = זוג SM, חשבו כמה TPC-ים מופעלים בכרטיס.
2. פתחו את ה-whitepaper של הארכיטקטורה שלכם (למשל "NVIDIA H100 Tensor Core GPU Architecture") ומצאו את מספר ה-SM-ים של השבב המלא (על GH100: 144) ואת מספר ה-GPC-ים (על GH100: 8).
3. חשבו את ההפרש בין השבב המלא לכרטיס שברשותכם, והסבירו במילה אחת למה ההפרש קיים.
4. ענו: מדוע התוכנית מתרגיל 1 יכולה לתת לכם את מספר ה-SM אבל לא את מספר ה-GPC-ים? מה זה אומר על מה שה-runtime של CUDA בכלל חושף?

רמז: על H100, 132 / 2 = 66 TPC-ים מופעלים. ההפרש 144 - 132 = 12 SM-ים כבויים נובע מתשואת ייצור (SM-ים פגומים שכובו). אין שדה gpcCount ב-cudaDeviceProp - ה-runtime חושף multiProcessorCount אבל לא את חלוקת ה-GPC/TPC, ולכן את המספרים האלה קוראים רק מה-whitepaper.

תרגיל 3 - מיפוי המונחים לתרשים

1. ציירו (על נייר או ב-ASCII) את שרשרת ההכלה מ-GPU ועד ליבה בודדת, עם כל חמשת הרבדים: GPU, GPC, TPC, SM, ליבה.
2. סמנו על התרשים את כל הכמויות שאתם בטוחים בהן מהתרגילים הקודמים (מספר SM, מספר TPC), ובנפרד את אלה שקראתם מה-whitepaper (מספר GPC, TPC ל-GPC).
3. הוסיפו לתרשים את שני הרכיבים ה"מיוחדים" שאינם SM: היכן יושב ה-raster engine, ומה מגדיר TPC מלבד היותו זוג.
4. סמנו בחץ: איפה משובץ block בודד, ואיפה (ב-Hopper) משובץ cluster של block-ים.

רמז: ה-raster engine יושב ברמת ה-GPC (אחד לכל GPC). ההגדרה של TPC היא "זוג SM צמודים פיזית" - הצמידות היא חלק מההגדרה, לא פרט טכני. block משובץ ל-SM יחיד; cluster משובץ ל-GPC יחיד. אם התרשים שלכם נכון, המבנה GPU → GPC → TPC → SM → core צריך להיות עץ שבו לכל TPC בדיוק שני עלי-SM.

תרגיל 4 - למה זיכרון משותף בקלאסטר, אבל לא בין block-ים שרירותיים

1. נסחו במילים שלכם את הכלל: block משובץ ל-SM יחיד, ו-cluster (ב-Hopper) משובץ ל-GPC יחיד.
2. הסבירו, על בסיס שרשרת ההכלה בלבד, מדוע שני block-ים באותו cluster יכולים לגשת ל-shared memory זה של זה (distributed shared memory).
3. הסבירו מדוע שני block-ים שרירותיים ב-grid (שאינם באותו cluster) אינם יכולים לעשות זאת.
4. הכלילו: מהי התכונה הפיזית שהופכת גישה מבוזרת לזיכרון לאפשרית, ולמה היא נעצרת בגבול ה-GPC?

רמז: shared memory יושב פיזית בתוך ה-SM. שני block-ים באותו cluster רצים על SM-ים באותו GPC, ולכן החומרה יכולה לנתב גישה ביניהם. שני block-ים שרירותיים עלולים לרוץ על GPC-ים שונים לגמרי - אין ביניהם ערובת מיקום, ולכן אין מסלול חומרתי לזיכרון של השני. הגבול הוא ה-GPC כי מעבר לו אין עוד הבטחת מיקום פיזי.

תרגיל 5 - השוואת דורות: A100 מול H100 מול B200

1. בנו טבלה עם עמודות: כרטיס, ארכיטקטורה, compute capability, מספר SM מופעל, מספר SM בשבב מלא.
2. מלאו את A100 SXM ו-H100 SXM (השתמשו במספרים מההרצאה: 108/128 ו-132/144).
3. חקרו את B200: מצאו ב-whitepaper של Blackwell או במפרט הרשמי כמה SM-ים בו, וציינו את התכונה המבנית שמייחדת אותו מהשניים האחרים.
4. ענו: איזה מספר נשאר קבוע לגמרי על פני כל שלושת הכרטיסים, ומדוע?

רמז: A100 (Ampere, CC 8.0), H100 (Hopper, CC 9.0), B200 (Blackwell, CC 10.0). ה-B200 בנוי משני dies הפועלים כ-GPU לוגי אחד - שינוי מבני משמעותי לעומת שבב יחיד. הקבוע היחיד: TPC = 2 SM, בכל הדורות בלי יוצא מן הכלל. אם אתם מריצים על אחד מהם, הצליבו את multiProcessorCount מול ה-whitepaper.

תרגיל 6 (בונוס) - איתור ה-.cta_group ב-SASS של Blackwell

1. אם יש לכם גישה לכרטיס Blackwell (או ל-CUDA Toolkit עדכני), כתבו kernel קטן שמשתמש ב-Tensor Cores דרך ספריית cuda::wmma או דרך CUTLASS, והדרו אותו ל--arch=sm_100a.
2. פרקו את הבינארי עם cuobjdump -sass ./a.out וחפשו פקודות MMA. שימו לב אם מופיע הסימון 1SM או 2SM.
3. הסבירו את הקשר: איזה מְאַפְיֵן PTX (.cta_group::1 או .cta_group::2) מוליד כל אחת מהגרסאות, ולאיזה רובד חומרה כל אחת מכוונת.
4. אם אין לכם כרטיס Blackwell: קראו בתיעוד ה-PTX ISA של NVIDIA את הערך המדויק של tcgen05, ותארו במילים איך שדה ה-.cta_group הוא "הפעם הראשונה שה-TPC ניתן למיעון בקוד".

רמז: .cta_group::1 → SASS מסוג 1SM (SM בודד); .cta_group::2 → SASS מסוג 2SM (זוג ה-SM-ים של ה-TPC). זו החלטת קומפילציה שמשנה התנהגות בזמן ריצה. cuobjdump -sass מדפיס את קוד המכונה של ה-GPU; אם אינכם רואים MMA כלל, ודאו שהקומפילציה כללה שימוש ב-Tensor Cores ולא רק ב-FP32 רגיל. הצמידות הפיזית של שני ה-SM-ים ב-TPC היא מה שהופך את וריאנט ה-2SM לאפשרי מלכתחילה.