2.5 Compute Capability וגרסאות ארכיטקטורה הרצאה
לאורך פרק 2 בנינו את מודל התכנות של CUDA מלמעלה למטה: בשיעור 2.1 פגשנו את ה-thread (thread), את ה-warp (קבוצת 32 threads) ואת מודל ה-SIMT (Single Instruction, Multiple Threads); בשיעור 2.2 ראינו איך block (בלוק) ו-grid (רשת) ממופים לחומרה; בשיעור 2.3 כתבנו kernel (קרנל, פונקציית ה-GPU) עם אינדוקס נכון ו-launch configuration; ובשיעור 2.4 סנכרנו threads עם barriers ו-cooperative groups. כל אלה הם הפשטות ניידות - הן לא מזכירות דגם כרטיס אחד. אבל בסופו של יום הקוד חייב לרוץ על שבב פיזי מסוים, וכאן נכנס המושג שסוגר את הפרק: ה-compute capability, מספר הגרסה שמחבר את הקוד שכתבנו לחומרה שתריץ אותו. בשיעור 1.2 כבר נגענו בו בקצרה כשדיברנו על ארכיטקטורת ה-SM; כאן נצלול אליו לעומק - נבין למה בכלל צריך מספר גרסה, מה ההבדל בין ארכיטקטורה וירטואלית לפיזית, איך עובד "מודל הבצל" של תאימות קדימה, מה שוברות הסיומות a ו-f החדשות, ואיך קוראים את ה-compute capability של הכרטיס שלכם בשלוש דרכים. את הטיפול המלא בדגלי הcompilation של nvcc (CUDA compiler) נשאיר לפרק 4, אבל התשתית המושגית נבנית כאן.
הבעיה ש-Compute Capability פותרת - Versioning the ISA¶
בשיעורי פרק 4 נראה שכל kernel עובר בסוף שרשרת compilation לשתי שפות נמוכות: PTX (Parallel Thread Execution) - אסמבלי וירטואלי נייד, ו-SASS (Streaming Assembler) - קוד המכונה הילידי שה-SM (מעבד רב-זרמי) מריץ בפועל. הבעיה: פקודות ב-PTX תואמות רק לחלק מהכרטיסים הפיזיים. פקודת Tensor Core (ליבת טנזור) שקיימת ב-Hopper לא קיימת ב-Pascal; גודל ה-warp, מספר האוגרים, סוגי הפעולות - כולם משתנים בין דורות. אם הcompiler היה צריך להכיר כל שבב ושבב בשמו, כל שחרור כרטיס חדש היה שובר את כל שרשרת הכלים.
הפתרון הוא שכבת הפשטה: מערכת מספור גרסאות שמפרידה את פרטי החומרה הפיזית מקבוצת הפקודות ומהcompiler. המערכת הזו נקראת compute capability. במקום לכוון לשבב ספציפי, הקוד מכוון למשפחה של כרטיסים בעלי אותה יכולת. זו בדיוק ההגדרה מהמקור: "מערכת המספור שמשמשת להפשטת פרטי ה-GPU הפיזי מקבוצת הפקודות ומהcompiler נקראת Compute Capability". ה-compute capability הוא גבול ההפשטה בין שני העולמות:
code you write abstraction boundary physical hardware
+-------------------+ +----------------------+ +-------------+
| CUDA C++ kernel | | | | H100 9.0 |
| | | -----> | Compute Capability | ----> | A100 8.0 |
| v | | major.minor | | RTX 4090 |
| PTX / SASS | | (7.0, 8.6, 9.0 ...) | | 8.9 |
+-------------------+ +----------------------+ +-------------+
"to which capability" not "to which chip" "which chips"
הרעיון המרכזי: מכוונים ליכולת, לא לשבב. kernel שקומפל ליכולת 8.0 ירוץ על כל כרטיס Ampere בעל היכולת הזו - A100, A30, A10 - בלי לדעת דבר על הדגם המסוים. זו אותה תכונת מדרגיות שראינו ב-grid בשיעור 2.2, אבל במימד הזמן: לא רק "אותו קוד על GPU קטן וגדול", אלא "אותו קוד על דורות חומרה שונים".
ארכיטקטורה וירטואלית מול פיזית - compute_XX מול sm_XX¶
הנקודה שהכי מבלבלת מתחילים היא שיש שני מספרי גרסה, לא אחד, והם נראים כמעט זהים. תיעוד nvcc קורא ל-compute capability "ארכיטקטורת GPU וירטואלית" (virtual GPU architecture), בניגוד ל"ארכיטקטורת GPU פיזית" (physical GPU architecture) שמבוטאת על ידי גרסת ה-SM. שני העולמות מקבלים תחביר משלהם:
- ארכיטקטורה וירטואלית -
compute_XX: היעד של ה-PTX. זו רמת היכולת המופשטת.compute_90אומר "PTX שמניח את היכולות של Hopper", בלי להתחייב לשבב. - ארכיטקטורה פיזית -
sm_XX: היעד של ה-SASS. זהו קוד המכונה הילידי לגרסת SM ספציפית.sm_90אומר "SASS שרץ ישירות על SM בגרסה 9.0".
| היבט | ארכיטקטורה וירטואלית | ארכיטקטורה פיזית |
|---|---|---|
| תחביר | compute_XX |
sm_XX |
| מכוונת אליה | PTX | SASS |
| מהי | רמת יכולת מופשטת | קוד מכונה לגרסת SM |
| מתי מתורגמת לחומרה | בזמן טעינה (JIT) | כבר מוכן, רץ ישירות |
| ניידות קדימה | כן (רצה על דורות חדשים) | לא (קשורה לגרסה הראשית) |
היחס ביניהם: PTX (וירטואלי) הוא שלב ביניים; SASS (פיזי) הוא היעד הסופי. אפשר לכוון ל-PTX ולתת לדרייבר להשלים את הדרך ל-SASS בזמן ריצה, או לכוון ישירות ל-SASS ולחסוך את השלב הזה. התנהגות ברירת המחדל של nvcc: כשמכוונים ל-compute capability, הcompiler "גם מייצר SASS ממוטב עבור ארכיטקטורת ה-SM התואמת". כלומר, -arch=sm_90 איננו רק "SASS ל-H100" - הוא בונה SASS ל-sm_90 וגם מטמיע PTX של compute_90 באותו artifact, כדי שתהיה גם ניידות קדימה. נפרק את המשמעות המעשית של זה בסעיף האחרון, ואת התחביר המלא של -gencode בפרק 4.
nvcc -arch=sm_90 foo.cu
|
+--> compute_90 -> PTX (virtual, forward-portable) ---+
| |--> same fatbin
+--> sm_90 -> SASS (physical, for H100, fast) ----+
מודל הבצל - תאימות קדימה - The Onion Model¶
מה קורה כשמריצים artifact ישן על כרטיס חדש? כאן נכנס עקרון התאימות המרכזי. עבור גרסאות סטנדרטיות (בלי סיומת), NVIDIA "מבטיחה תאימות קדימה (קוד PTX ישן רץ על כרטיסים חדשים) גם בין גרסאות ראשיות וגם בין גרסאות משניות, לפי מודל שכבות הבצל". השם "בצל" קולע: כל גרסה חדשה עוטפת את הישנות כשכבה נוספת, וה-PTX של שכבה פנימית (ישנה) יכול לרוץ בכל שכבה חיצונית (חדשה) שעוטפת אותה.
The onion-layers model - PTX forward compatibility
+---------------------------------------------------+
| Blackwell 12.0 |
| +-------------------------------------------+ |
| | Hopper 9.0 | |
| | +-----------------------------------+ | |
| | | Ampere 8.0 | | |
| | | +---------------------------+ | | |
| | | | Volta 7.0 (old PTX) | | | |
| | | +---------------------------+ | | |
| | +-----------------------------------+ | |
| +-------------------------------------------+ |
+---------------------------------------------------+
PTX of compute_70 -> runs on 7.0, 8.0, 9.0, 12.0
the arrow always points from inside out: old runs on new, never the other way
המנגנון שמאפשר את זה הוא קומפילציה בזמן ריצה (JIT - Just-In-Time) בתוך הדרייבר. כשטוענים artifact שמכיל PTX של compute_70 על כרטיס H100 (9.0), הדרייבר מזהה שאין SASS מתאים, לוקח את ה-PTX ומקמפל אותו ל-SASS של sm_90 ברגע הטעינה. התוצאה נשמרת ב-cache על הדיסק (ברירת המחדל: ~/.nv/ComputeCache), כך שההרצה הבאה לא משלמת שוב את מחיר ה-JIT.
loading an artifact with PTX compute_70 on H100 (9.0)
------------------------------------------------
[ is there SASS for sm_90? ] --no--> [ driver compiles PTX -> SASS for sm_90 ]
|
v
[ saved in ~/.nv/ComputeCache ]
|
v
[ run. subsequent runs: from the cache ]
שימו לב לכיוון: תאימות קדימה, לא אחורה. PTX ישן רץ על כרטיס חדש (JIT מעלה אותו), אבל SASS או PTX שנבנו ליכולת חדשה לא ירוצו על כרטיס ישן - הפקודות פשוט לא קיימות שם. אין דרך "להנמיך" יכולת.
חריגות מהבצל - הסיומות a ו-f - The a and f Suffixes¶
הדורות האחרונים שברו את אחידות מודל הבצל בכוונה, כי חלק מהיכולות החדשות (בעיקר סביב ה-Tensor Cores) קשורות כל כך חזק לחומרה ספציפית עד שאי אפשר להבטיח עבורן ניידות קדימה. לכן נוספו שתי סיומות למספר הגרסה:
- סיומת
a- הוצגה עם Hopper (למשל9.0a): מסמנת יכולות שחורגות ממודל הבצל. עבורן, "תאימות עתידית אינה מובטחת, אפילו בתוך אותה גרסה ראשית". אלה היכולות ה"מואצות" (architecture-specific), כמו פקודות ה-wgmmaוה-TMA (Tensor Memory Accelerator) של Hopper. קוד שנבנה ל-sm_90aתפור ל-Hopper ואין שום ערובה שירוץ על דור עתידי כלשהו, גם לא על גרסה ראשית 9 אחרת. - סיומת
f- הוצגה עם Blackwell (למשל10.0f): גם היא חורגת ממודל הבצל, אבל "קרובה יותר ל-SemVer": התאימות מובטחת בין גרסאות משניות אבל לא בין גרסאות ראשיות. כלומרsm_100fיהיה תואם על פני משפחת 10.x, אבל לא יקפוץ ל-12.x ולא יירד ל-9.x.
הנה שלושת המשטרים זה מול זה:
| משטר | דוגמה | תאימות בין גרסאות משניות | תאימות בין גרסאות ראשיות | הוצג בדור |
|---|---|---|---|---|
| סטנדרטי (בצל) | sm_90, compute_80 |
כן (קדימה) | כן (קדימה) | מאז ומתמיד |
סיומת a |
sm_90a |
לא מובטח | לא מובטח | Hopper |
סיומת f |
sm_100f |
כן | לא | Blackwell |
הלקח המעשי: אי אפשר יותר להניח "PTX ישן תמיד רץ בכל מקום". ברגע שאתם משתמשים ביכולת מואצת של דור מסוים (Tensor Cores מהדור החדש, TMA), אתם עוברים לתחום a או f ומאבדים חלק מהערובות של מודל הבצל. צריך להכיר את כללי הגרסה הספציפיים ולא לסמוך על אחידות שכבר לא קיימת. בפרק 4 נראה מתי nvcc מכריח אתכם לסיומת a (למשל, פקודות wgmma פשוט לא מתקמפלות תחת sm_90 רגיל - רק תחת sm_90a).
מפת ה-Compute Capability לארכיטקטורות - The CC-to-Architecture Map¶
עד כאן דיברנו על המבנה של המספר. עכשיו נחבר מספרים לשמות. אזהרה חשובה: עמוד המקור של ה-compute capability אינו מונה את המפה הזו - הוא מפנה למדריך התכנות של NVIDIA לטבלאות היכולות המלאות. המפה שלהלן היא ידע CUDA סטנדרטי שאנחנו מוסיפים לנוחותכם כתוספת חיצונית שימושית, לא ציטוט מהמקור:
| Compute Capability | משפחת ארכיטקטורה | כרטיס לדוגמה |
|---|---|---|
| 7.0 / 7.2 | Volta | V100 |
| 7.5 | Turing | T4, RTX 2080 |
| 8.0 | Ampere | A100 |
| 8.6 / 8.7 | Ampere | RTX 3090, Orin |
| 8.9 | Ada Lovelace | RTX 4090, L4 |
| 9.0 | Hopper | H100, H200 |
| 10.x / 12.x | Blackwell | B200, RTX 5090 |
המלכודת החשובה ביותר בטבלה - שורות ה-8.x: הגרסה הראשית כמעט נרדפת למשפחת הארכיטקטורה (כל 7.x = Volta/Turing, כל 9.x = Hopper), אבל "כמעט" איננו "תמיד". Ampere משתמשת ב-8.0, 8.6, 8.7, ואילו Ada Lovelace משתמשת ב-8.9 - אותה גרסה ראשית (8) כמו Ampere, אבל משפחת ארכיטקטורה שונה לחלוטין. גרסה ראשית משותפת אינה מבטיחה אותה משפחה. אם תסתמכו על "major==8 אז זה Ampere", תיכשלו על כל RTX 4090. תמיד בדקו את המספר המלא major.minor, לא רק את הראשי.
מהמפה הזו נגזרת גם דוגמת הריצה של הקורס: ה-H100 SXM שלנו הוא compute capability 9.0, משפחת Hopper. כשנכתוב sm_90 נכוון בדיוק אליו, וכשנרצה את היכולות המואצות של Hopper (Tensor Cores מהדור הרביעי, TMA) נכתוב sm_90a.
איך לקרוא את ה-Compute Capability של הכרטיס - Reading Your CC¶
יש שלוש דרכים לגלות את ה-compute capability, לפי ההקשר שבו אתם נמצאים. כדאי להכיר את שלושתן.
דרך 1 - משורת הפקודה עם nvidia-smi. המהירה ביותר, בלי לכתוב שורת קוד:
על H100 נקבל:
דרך 2 - בזמן ריצה, מקוד ה-host (המארח) עם cudaGetDeviceProperties. כך תבחרו התנהגות דינמית לפי הכרטיס. שימו לב למקרו CUDA_CHECK שעוטף כל קריאת CUDA ומדפיס קובץ, שורה ושגיאה אם משהו נכשל - נשתמש בו בכל תוכנית host לאורך הקורס:
#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) {
int dev = 0;
cudaDeviceProp p;
CUDA_CHECK(cudaGetDeviceProperties(&p, dev));
printf("GPU : %s\n", p.name);
printf("Compute Capability : %d.%d\n", p.major, p.minor);
printf("SM arch target : sm_%d%d\n", p.major, p.minor);
return 0;
}
הcompilation והרצה (החליפו את הארכיטקטורה בזו של הכרטיס שלכם):
הפלט על H100:
השדות p.major ו-p.minor הם בדיוק שני חלקי ה-compute capability. זו הדרך המהימנה ביותר, כי היא שואלת את הכרטיס עצמו במקום לשנן.
דרך 3 - בזמן compilation, מקוד ה-device (ההתקן) עם __CUDA_ARCH__. בתוך kernel אפשר לבחור מסלול קוד שונה לכל יכולת. המקרו __CUDA_ARCH__ מוגדר רק במעברי הcompilation של ה-device (בקוד ה-host הוא לא מוגדר), וערכו הוא major*100 + minor*10. עבור sm_90 הוא 900, עבור sm_80 הוא 800:
__global__ void k(float* out) {
#if __CUDA_ARCH__ >= 900
// Hopper and above path
out[threadIdx.x] = 9.0f;
#elif __CUDA_ARCH__ >= 800
// Ampere path
out[threadIdx.x] = 8.0f;
#else
// common path for older generations
out[threadIdx.x] = 0.0f;
#endif
}
עדינות שקשורה לסיומת a: כשמקמפלים ל-sm_90a, הערך של __CUDA_ARCH__ נשאר 900 (זהה ל-sm_90 הרגיל), אבל nvcc מגדיר בנוסף מקרו נפרד, __CUDA_ARCH_FEAT_SM90_ALL, שמסמן שהיכולות המואצות זמינות. כך קוד שמשתמש ב-wgmma יכול לשמור על עצמו מפני compilation תחת sm_90 רגיל, שבו הפקודות האלה כלל לא קיימות.
מה זה אומר בפועל - PTX מול SASS בזמן טעינה - Ship PTX or SASS¶
נסגור עם ההשלכה המעשית שמנחה כל החלטת compilation, ושנעמיק בה בפרק 4 עם הטיפול המלא ב--gencode. יש כאן פשרה (trade-off) חדה בין שתי גישות:
גישה א - לשלוח PTX (ניידות קדימה). אם ה-artifact מכיל PTX של compute_XX, הוא ירוץ על כל דור עתידי בזכות ה-JIT של הדרייבר. המחיר: בטעינה הראשונה על כל כרטיס חדש, הדרייבר עוצר ומקמפל את ה-PTX ל-SASS - זמן ה-JIT. עבור אפליקציה עם עשרות kernels זה יכול להוסיף שניות להפעלה הראשונה. בנוסף, ה-SASS שה-JIT מייצר אינו בהכרח ממוטב כמו SASS שנבנה ידנית לאותו sm, כי ה-JIT עובד תחת אילוצי זמן ומידע.
גישה ב - לכוון ל-sm_XX המדויק (ביצועים). אם ה-artifact מכיל SASS מוכן ל-sm_90, הטעינה על H100 מיידית - אין JIT, וה-SASS ממוטב היטב. המחיר: אין ניידות קדימה. אם אין PTX בתמונה, הרצה על כל גרסה ראשית אחרת פשוט תיכשל בטעינה.
the two ends of the trade-off
-----------------------------------------------------------
PTX only (compute_90) | SASS only (sm_90) | both (default -arch=sm_90)
---------------------|--------------------|----------------------------
forward portability: yes | portability: no | portability + speed (larger artifact)
load cost: JIT | load cost: 0 | 0 on sm_90, JIT on newer generations
performance: medium | performance: optimal | optimal on sm_90
בגלל זה ברירת המחדל של nvcc מטמיעה את שניהם כשכותבים -arch=sm_90: SASS ל-sm_90 (מהיר ומיידי על ה-H100 שלכם עכשיו) ו-PTX של compute_90 (שיציל אתכם על כרטיס עתידי דרך JIT). זו הפשרה הסבירה כברירת מחדל. בפרק 4 נראה איך -gencode נותן שליטה מדויקת - להטמיע SASS למספר sm-ים בו-זמנית (fatbin), להשמיט את ה-PTX כשרוצים artifact קטן, או להוסיף PTX של יכולת גבוהה כרשת ביטחון קדימה. אם תרצו לכפות JIT לצורך בדיקה, אפשר עם משתנה הסביבה CUDA_FORCE_PTX_JIT=1, ולנקות את ה-cache עם CUDA_CACHE_DISABLE=1.
סיכום¶
- ה-compute capability הוא מערכת מספור גרסאות שמפשיטה את פרטי ה-GPU הפיזי מקבוצת הפקודות (PTX) ומהcompiler, כך שהקוד מכוון למשפחת כרטיסים בעלת יכולת נתונה, לא לשבב מסוים.
- הפורמט הוא
major.minor(למשל 7.0, 8.6, 9.0); יש להבחין בין ארכיטקטורה וירטואלית (compute_XX, היעד של ה-PTX) לארכיטקטורה פיזית (sm_XX, היעד של ה-SASS), וברירת המחדל של nvcc מייצרת גם SASS ממוטב ל-SM התואם. - מודל שכבות הבצל מבטיח לגרסאות סטנדרטיות תאימות קדימה: PTX ישן רץ על כרטיסים חדשים יותר, בין גרסאות ראשיות ומשניות, באמצעות JIT בזמן טעינה בתוך הדרייבר, עם cache על הדיסק ב-
~/.nv/ComputeCache. - הכיוון הוא תאימות קדימה בלבד - PTX ישן עולה לחומרה חדשה, אבל קוד שנבנה ליכולת חדשה לעולם לא ירוץ על כרטיס ישן; אין הנמכת יכולת.
- הסיומת
a(Hopper, למשל9.0a) מסמנת יכולות מואצות ששוברות את מודל הבצל ללא שום ערובת תאימות עתידית, אפילו בתוך אותה גרסה ראשית; הסיומתf(Blackwell, למשל10.0f) קרובה יותר ל-SemVer - תואמת בין גרסאות משניות אך לא בין ראשיות. - מפת היכולות לארכיטקטורות (Volta 7.0, Turing 7.5, Ampere 8.0/8.6, Ada 8.9, Hopper 9.0, Blackwell 10.x/12.x) היא תוספת חיצונית שאינה מופיעה בעמוד המקור; המלכודת המרכזית בה היא ש-Ada (8.9) חולקת גרסה ראשית 8 עם Ampere אך היא משפחה נפרדת, ולכן חובה לבדוק
major.minorמלא. - קוראים את ה-compute capability בשלוש דרכים:
nvidia-smi --query-gpu=compute_capמשורת הפקודה,cudaGetDeviceProperties(השדותmajor/minor) בקוד ה-host עטוף ב-CUDA_CHECK, ו-__CUDA_ARCH__(ערךmajor*100+minor*10) בקוד ה-device בזמן compilation. - הפשרה המעשית: שליחת PTX נותנת ניידות קדימה אך גובה מחיר JIT בטעינה ו-SASS פחות ממוטב, בעוד כיוון ל-
sm_XXהמדויק נותן SASS מיטבי וטעינה מיידית אך ללא ניידות; את השליטה המלאה עם-gencodeותהליך ה-fatbin נלמד בפרק 4.