1.1 מבנה ה GPU GPC, TPC ו SM פתרון
פתרון - מבנה ה-GPU - GPC, TPC ו-SM¶
כאן נעבור על כל תרגיל עם הפקודות, הקוד והפלט המדויקים. שימו לב: כל המספרים בהמשך הם מהרצה על H100 SXM, והם ישתנו בין כרטיסים - זו כל המטרה של תרגילי החקירה האלה. המבנה, השיטה וההיגיון זהים בכל מכונה. אם אתם על T4 ב-Colab, החליפו -arch=sm_90a ב--arch=sm_75 ותקבלו מספרים אחרים (ל-T4 יש 40 SM-ים).
פתרון תרגיל 1 - שלוש דרכים לספור SM-ים¶
שלב 1 - nvidia-smi. הפלט (מקוצר):
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 550.xx Driver Version: 550.xx CUDA Version: 12.4 |
| 0 NVIDIA H100 80GB HBM3 ... 81559MiB ... |
+-----------------------------------------------------------------------------+
שם הכרטיס וכ-80 GiB זיכרון - אבל אין שום שורה עם מספר SM. זו ההפתעה: הכלי הכי מוכר לא נותן את המספר שאנחנו רוצים.
שלב 2 - התוכנית שלנו. הקובץ smcount.cu:
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d\n '%s'\n -> %s\n", \
__FILE__, __LINE__, #call, cudaGetErrorString(err_)); \
exit(EXIT_FAILURE); \
} \
} while (0)
int main(void) {
cudaDeviceProp p;
CUDA_CHECK(cudaGetDeviceProperties(&p, 0));
printf("%s\n", p.name);
printf("compute capability: %d.%d\n", p.major, p.minor);
printf("multiProcessorCount (SMs): %d\n", p.multiProcessorCount);
printf("implied CUDA cores (FP32): %d\n", p.multiProcessorCount * 128);
return 0;
}
הcompilation והרצה:
הפלט:
NVIDIA H100 80GB HBM3
compute capability: 9.0
multiProcessorCount (SMs): 132
implied CUDA cores (FP32): 16896
שלב 3 - deviceQuery. השורה הרלוונטית מתוך הפלט הארוך:
שלב 4 - עקביות. שלושת המקורות מסכימים על 132 SM-ים. מה שחשוב לרשום: nvidia-smi חשף שם, נהג, גרסת CUDA וזיכרון - אבל לא מספר SM; המקור למספר ה-SM הוא multiProcessorCount (דרך התוכנית שלנו או דרך deviceQuery).
למה זה עבד: cudaGetDeviceProperties ממלא מבנה cudaDeviceProp מהמנהל, ובו השדה multiProcessorCount הוא בדיוק מספר ה-SM-ים הפעילים בכרטיס. הכפל ב-128 עובד כי לכל SM ב-H100 יש 128 יחידות CUDA Core מסוג FP32, ולכן 132 × 128 = 16896 - אותו מספר הthreads המקביליים-באמת שחישבנו בשיעור 0.1.
איך להכליל: cudaDeviceProp הוא הדרך התוכנתית הקנונית לחקור כל כרטיס בזמן ריצה - ממנו שולפים גם sharedMemPerBlock, warpSize, maxThreadsPerBlock ועוד. nvidia-smi נועד לניטור וניהול (טמפרטורה, הספק, תהליכים), לא לארכיטקטורה - אל תחפשו בו מבנה חומרתי. תמיד קראו את multiProcessorCount מהכרטיס במקום לקודד מספר קשיח בתוכנה שאמורה לרוץ על כמה כרטיסים.
פתרון תרגיל 2 - שחזור ה-TPC וההצלבה מול ה-whitepaper¶
שלב 1 - מ-SM ל-TPC. מכיוון ש-TPC = זוג SM צמודים:
שלב 2 - השבב המלא. מה-whitepaper של GH100:
שלב 3 - ההפרש. ההפרש בין השבב המלא לכרטיס שלנו:
המילה האחת שמסבירה זאת: תשואה (yield). בשבב ענק כמעט תמיד יש פגמים נקודתיים; NVIDIA מכבה את ה-SM-ים הפגומים ומוכרת שבב עם 132 SM תקינים במקום 144.
שלב 4 - GPC מול SM. התוכנית מתרגיל 1 נותנת SM אבל לא GPC כי ה-runtime של CUDA פשוט אינו חושף את חלוקת ה-GPC/TPC - אין שדה gpcCount או tpcCount ב-cudaDeviceProp. מה שה-runtime חושף הוא מה שרלוונטי למודל התכנות: כמה SM-ים יש (כי block משובץ ל-SM), כמה זיכרון משותף, גודל warp וכו'. מבנה ה-GPC היה, עד Hopper, פרט חומרה בלבד - ולכן הוא חי ב-whitepaper, לא ב-API.
למה זה עבד: ההגדרה "TPC = 2 SM" היא קבוע ארכיטקטורי, ולכן חלוקת מספר ה-SM ב-2 נותנת ישירות את מספר ה-TPC-ים המופעלים. ה-whitepaper הוא המקום היחיד שבו מפורטת חלוקת ה-8 GPC × 9 TPC, כי זו החלטה של תכנון השבב שלא באה לידי ביטוי ב-API של זמן ריצה.
איך להכליל: לכל דור יש whitepaper עם "full-chip diagram" שממנו קוראים GPC, TPC ל-GPC, ו-SM בשבב מלא; את המספר המופעל בפועל קוראים מהכרטיס. הכלל: מבנה (כמה GPC/TPC) מגיע מה-whitepaper, כמות מופעלת (כמה SM) מגיעה מהכרטיס. ההפרש ביניהם הוא תמיד סיפור של תשואת ייצור.
פתרון תרגיל 3 - מיפוי המונחים לתרשים¶
התרשים המלא, עם הכמויות של H100 מסומנות:
GPU: NVIDIA H100 SXM
│ ← 80 GiB HBM3
├── GPC 0 [raster engine] ← 8 GPCs in the full GH100 chip
│ ├── TPC 0 (a pair of adjacent SMs) ← up to 9 TPC per GPC ; 66 TPC enabled in total
│ │ ├── SM 0 ─┐
│ │ └── SM 1 ─┴─ a block is scheduled onto a single SM
│ ├── TPC 1
│ │ ├── SM 2
│ │ └── SM 3
│ └── ...
├── GPC 1 [raster engine] ─┐
├── ... ├─ a cluster (Hopper) is scheduled onto a single GPC
└── GPC 7 [raster engine] ─┘
← 132 SM enabled ; 132 × 128 = 16896 CUDA Cores
inside a single SM:
4 × Warp Scheduler · 128 × CUDA Core (FP32) · 4 × Tensor Core · 256 KiB L1/shared
התשובות המשולבות בתרשים: ה-raster engine יושב ברמת ה-GPC (אחד לכל GPC). ה-TPC מוגדר כזוג SM צמודים פיזית - הצמידות היא חלק מההגדרה. block משובץ ל-SM יחיד; cluster משובץ ל-GPC יחיד.
למה זה עבד: העץ נובע ישירות משרשרת ההכלה GPU → GPC → TPC → SM → core. כל TPC הוא בדיוק זוג עלי-SM, ולכן מספר ה-SM תמיד זוגי ומתחלק ב-2 למספר TPC. הפרדת ה"בטוח" (SM, TPC - מחושבים מהכרטיס) מה"נקרא מ-whitepaper" (GPC, TPC ל-GPC) משקפת מה ה-runtime באמת חושף.
איך להכליל: אותו עץ מתאים לכל כרטיס NVIDIA מודרני - רק המספרים משתנים. כשאתם פוגשים כרטיס חדש, מלאו את העץ הזה: SM מ-multiProcessorCount, TPC = SM/2, GPC מה-whitepaper. זהו כלי אבחון מהיר להבנת כל שבב.
פתרון תרגיל 4 - למה זיכרון משותף בקלאסטר, אבל לא בין block-ים שרירותיים¶
שלב 1 - הכלל. block משובץ ל-SM יחיד (החומרה מבטיחה שכל threads שלו על אותו SM); cluster ב-Hopper משובץ ל-GPC יחיד (החומרה מבטיחה שכל ה-block-ים שלו על SM-ים באותו GPC).
שלב 2 - למה כן בקלאסטר. ה-shared memory יושב פיזית בתוך ה-SM. שני block-ים באותו cluster רצים על שני SM-ים שנמצאים באותו GPC. מכיוון שיש הבטחת מיקום כזו, החומרה יכולה לספק מסלול ניתוב ישיר בין ה-SM-ים באותו GPC, וכך block אחד יכול לקרוא ולכתוב ל-shared memory של רעהו - זהו distributed shared memory.
שלב 3 - למה לא בין שרירותיים. לשני block-ים שרירותיים ב-grid אין שום הבטחת מיקום. המתזמן עלול לשבץ אותם על SM-ים בשני GPC-ים שונים לגמרי (או אפילו לא להריץ אותם בו-זמנית). בלי ערובה ששניהם על אותו GPC, אין מסלול חומרתי שמחבר את ה-shared memory של האחד לשני - ולכן גישה כזו פשוט אינה קיימת במודל.
שלב 4 - ההכללה. התכונה הפיזית שמאפשרת גישה מבוזרת היא קִרבת מיקום מובטחת: החומרה יודעת בוודאות ששני ה-block-ים על SM-ים סמוכים באותו GPC. הגבול נעצר ב-GPC כי זהו הרובד הגדול ביותר שבתוכו החומרה עדיין מבטיחה מיקום משותף; מעבר לגבול ה-GPC אין עוד ערובה, ולכן אין עוד גישה ישירה.
למה זה עבד: ההסבר כולו נגזר משרשרת ההכלה בלבד - "shared memory בתוך SM" + "cluster בתוך GPC" מחייבים שגישה מבוזרת תעבוד רק בגבולות ה-GPC. אין צורך בשום עובדה נוספת מעבר להיררכיה.
איך להכליל: זהו העיקרון הכללי בכל היררכיית זיכרון מקבילית: היקף השיתוף נקבע על ידי רובד המיקום המשותף. הthreads ב-block חולקים shared memory (רובד SM); block-ים ב-cluster חולקים distributed shared memory (רובד GPC); כל השאר חולק רק global memory (רובד ה-GPU כולו). כשתפגשו מנגנון סנכרון או שיתוף חדש, שאלו תמיד: "מהו הרובד הפיזי שבו המיקום המשותף מובטח?"
פתרון תרגיל 5 - השוואת דורות: A100 מול H100 מול B200¶
הטבלה המלאה:
Card Architecture CC SM enabled SM in full chip Structure
A100 SXM Ampere 8.0 108 128 8 GPC × up to 8 TPC × 2 SM
H100 SXM Hopper 9.0 132 144 8 GPC × up to 9 TPC × 2 SM
B200 Blackwell 10.0 ~160 (two dies) depends on SKU two dies as one logical GPU
שלב 4 - הקבוע. המספר שנשאר קבוע לחלוטין על פני שלושת הכרטיסים הוא 2 SM ל-TPC. מספר ה-SM, מספר ה-TPC ל-GPC, ואפילו מספר ה-dies - כולם משתנים; אבל ה-TPC נשאר זוג SM בכל דור בלי יוצא מן הכלל.
למה זה עבד: כל השורות נגזרות מאותה שרשרת הכלה GPU → GPC → TPC → SM, ורק הכמויות זזות בין הדורות. ה-B200 מדגים שאפילו קפיצה מבנית גדולה (מעבר משבב יחיד לשני dies) לא משנה את יחידת הבנייה הבסיסית - ה-TPC.
איך להכליל: כשמנתחים כרטיס חדש, מזהים קודם את הקבוע (TPC = 2 SM, ומכאן TPC = SM/2) ואז קוראים את המשתנים מה-whitepaper ומ-deviceQuery. B200 (שני dies) הוא תזכורת שהמונח "GPU" הוא לוגי - הוא יכול להתפרש על יותר מ-die פיזי אחד, אבל היררכיית ה-GPC/TPC/SM נשארת המסגרת שדרכה חושבים על כולם.
פתרון תרגיל 6 (בונוס) - איתור ה-.cta_group ב-SASS של Blackwell¶
שלב 1-2 - פירוק ה-SASS. לאחר compilation kernel שמשתמש ב-Tensor Cores ל--arch=sm_100a ופירוק עם cuobjdump:
פלט אופייני (מקוצר) יראה פקודות MMA מסומנות בווריאנט המתאים:
הסימון 1SM מול 2SM הוא הביטוי ב-SASS של שדה ה-.cta_group ברמת ה-PTX.
שלב 3 - הקשר PTX ל-SASS. ההתאמה:
PTX tcgen05.mma.cta_group::1 → SASS ...1SM → a single SM
PTX tcgen05.mma.cta_group::2 → SASS ...2SM → the TPC's SM pair
שלב 4 - למה זו הפעם הראשונה. לפני Blackwell, ה-TPC היה מבנה חומרה בלבד - שום פקודה, שום מְאַפְיֵן ושום API לא התייחסו אליו. שדה ה-.cta_group ב-tcgen05 הוא הראשון שמאפשר לקוד לבחור במפורש בין "הרץ על SM בודד" לבין "גייס את שני ה-SM-ים של ה-TPC יחד" - כלומר, לראשונה, ה-TPC ניתן למיעון בקוד.
למה זה עבד: cuobjdump -sass מדפיס את קוד המכונה הסופי (SASS) שאליו הcompiler תרגם את ה-PTX. מכיוון שבחירת .cta_group היא החלטת קומפילציה, היא מוטבעת כבר ב-SASS כשתי גרסאות פקודה נפרדות (1SM מול 2SM) - ולכן אפשר לראות אותה ישירות בפירוק. הצמידות הפיזית של שני ה-SM-ים ב-TPC היא מה שמאפשר לווריאנט 2SM לחלוק אופרנדים במסלולים קצרים ולעבוד יחד על אותה מטריצה.
איך להכליל: קריאת SASS עם cuobjdump -sass (ולפעמים nvdisasm) היא הדרך לאמת מה הcompiler באמת ייצר, מעבר למה שכתבנו ב-PTX או ב-CUDA C++ - נחזור לזה לעומק בפרק ה-PTX/SASS. הכלל הכללי: כשתכונת חומרה נעשית גלויה למתכנת, היא כמעט תמיד מופיעה תחילה כמְאַפְיֵן ב-PTX ואז כווריאנט פקודה ב-SASS; חיפוש הווריאנטים בפירוק הוא דרך אמינה לזהות אילו רבדים של החומרה הקוד שלכם באמת מנצל.