0.2 סביבת עבודה פתרון
פתרון - סביבת עבודה¶
נעבור תרגיל אחרי תרגיל עם הפקודות האמיתיות והפלט הצפוי. כל הפלטים כאן הם מ-H100 (SXM) עם דרייבר 550 ו-CUDA 12.4; אצלכם המספרים המדויקים - טמפרטורה, הספק, גרסאות, נתיבים - יהיו שונים, וזה תקין לגמרי. מי שעובד על T4 ב-Colab יראה מספרים אחרים לגמרי, ונציין אותם בתרגיל הבונוס.
פתרון תרגיל 1 - קריאת מצב הכרטיס¶
$ nvidia-smi
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 550.54.15 Driver Version: 550.54.15 CUDA Version: 12.4 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
|=========================================+========================+======================|
| 0 NVIDIA H100 80GB HBM3 On | 00000000:53:00.0 Off | 0 |
| N/A 32C P0 71W / 700W | 0MiB / 81559MiB | 0% Default |
+-----------------------------------------+------------------------+----------------------+
הקריאה של השדות:
- Driver Version 550.54.15 מול CUDA Version 12.4: הראשון הוא גרסת הדרייבר המותקן, השני הוא גרסת ה-CUDA המקסימלית שהדרייבר הזה תומך בה. הם שונים כי הם שני רכיבים שונים - הדרייבר וה-Toolkit מתעדכנים בנפרד. אפשר להתקין Toolkit 12.0 מול הדרייבר הזה, אבל לא Toolkit 12.6.
- Name: NVIDIA H100 80GB HBM3. Temp: 32C (בטל). Perf: P0 (מצב פעיל, תדר מלא). Pwr: 71W מתוך תקרה של 700W (ה-TDP של H100 SXM). Memory-Usage: 0MiB מתוך 81559MiB - הכרטיס ריק.
הקטעים הרלוונטיים מ-nvidia-smi -q:
Temperature
GPU Current Temp : 32 C
GPU Slowdown Temp : 85 C
GPU Shutdown Temp : 90 C
Clocks
SM : 345 MHz
Memory : 2619 MHz
סף ההאטה התרמי הוא 85C - מעליו הכרטיס יוריד תדר כדי להתקרר. שימו לב שתדר ה-SM בבטלה הוא 345 MHz בלבד, לעומת תדר ה-boost המלא של כ-1980 MHz תחת עומס - הכרטיס חוסך חשמל כשאין עבודה.
למה זה עבד: nvidia-smi הוא כלי מעל NVML ששואל את הדרייבר על מצב החומרה. עצם זה שהוא רץ מוכיח שהדרייבר טעון ותקין. איך להכליל: אותה קריאה מתאימה לכל כרטיס NVIDIA - רק המספרים משתנים. במכונה עם כמה כרטיסים כל אחד יקבל שורה משלו עם אינדקס נפרד, ותוכלו למקד כרטיס עם nvidia-smi -i 0.
פתרון תרגיל 2 - הוכחה שהדרייבר וה-Toolkit מותקנים¶
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Cuda compilation tools, release 12.4, V12.4.99
$ find / -name 'libcuda.so*' 2>/dev/null
/usr/lib/x86_64-linux-gnu/libcuda.so.1
/usr/lib/x86_64-linux-gnu/libcuda.so.550.54.15
$ find / -name 'libcudart.so*' 2>/dev/null
/usr/local/cuda-12.4/lib64/libcudart.so
/usr/local/cuda-12.4/lib64/libcudart.so.12
/usr/local/cuda-12.4/lib64/libcudart.so.12.4.99
שימו לב לנתיבים: libcuda.so יושבת תחת ספריות המערכת (/usr/lib/...) כי היא מגיעה מחבילת הדרייבר, ומספר הגרסה שלה (550.54.15) זהה לגרסת הדרייבר. libcudart.so יושבת תחת /usr/local/cuda-12.4/ כי היא מגיעה מה-Toolkit, ומספר הגרסה שלה (12.4) הוא גרסת ה-Toolkit.
התשובות: אם libcuda.so הייתה חסרה, זה היה אומר שהדרייבר לא מותקן או שבור - שום קוד CUDA לא היה רץ, גם לא nvidia-smi. אם רק nvcc היה חסר (בזמן ש-nvidia-smi עובד), היה לנו דרייבר בלי Toolkit - היינו יכולים להריץ תוכניות מקומפלות אבל לא לקמפל.
בדיקת הקישור של הבינארי אחרי הקימפול:
$ ldd ./devinfo
linux-vdso.so.1 (0x00007ffc...)
libcudart.so.12 => /usr/local/cuda-12.4/lib64/libcudart.so.12
libc.so.6 => /lib/x86_64-linux-gnu/libc.so.6
/lib64/ld-linux-x86-64.so.2
הבינארי נקשר דינמית ל-libcudart.so.12 (ה-Runtime API). שימו לב ש-libcuda.so דווקא לא מופיעה כאן - ה-Runtime טוען אותה בזמן ריצה דרך הדרייבר (dlopen), ולכן היא לא רשומה כתלות ישירה. זו נקודה עדינה: תוכנית Runtime API "רואה" רק את libcudart, וזו טוענת מתחתיה את הדרייבר.
למה זה עבד: כל חצי של המחסנית מגיע מחבילה נפרדת עם נתיב וגרסה משלו, ולכן אפשר לאמת כל אחד בנפרד. איך להכליל: ldd על כל בינארי CUDA יראה לכם מולן מ-הספריות הוא נקשר; אם libcudart לא מופיעה, הוא כנראה קושר סטטית (עם nvcc -cudart static), וזה נפוץ בבינארים מופצים.
פתרון תרגיל 3 - חקירת הכרטיס מתוך קוד¶
// devinfo.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: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while (0)
int main() {
int count = 0;
CUDA_CHECK(cudaGetDeviceCount(&count));
printf("CUDA devices found: %d\n\n", count);
for (int d = 0; d < count; ++d) {
cudaDeviceProp p;
CUDA_CHECK(cudaGetDeviceProperties(&p, d));
printf("Device %d: %s\n", d, p.name);
printf(" Compute capability: %d.%d\n", p.major, p.minor);
printf(" SM count: %d\n", p.multiProcessorCount);
printf(" Warp size: %d\n", p.warpSize);
printf(" Max threads / block: %d\n", p.maxThreadsPerBlock);
printf(" Shared mem / block: %zu KiB\n", p.sharedMemPerBlock / 1024);
printf(" Global (HBM) memory: %.1f GiB\n",
p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
printf(" GPU boost clock: %.0f MHz\n", p.clockRate / 1000.0);
}
return 0;
}
$ nvcc -arch=sm_90 devinfo.cu -o devinfo
$ ./devinfo
CUDA devices found: 1
Device 0: NVIDIA H100 80GB HBM3
Compute capability: 9.0
SM count: 132
Warp size: 32
Max threads / block: 1024
Shared mem / block: 48 KiB
Global (HBM) memory: 79.1 GiB
GPU boost clock: 1980 MHz
עבור H100: 132 SM, compute capability 9.0, כ-79 GiB HBM (מעט פחות מ-81559MiB של nvidia-smi, כי שם שמור לדרייבר), ו-warp של 32.
חישוב שיא ה-FP32:
בדיוק המספר הרשמי של H100.
למה זה עבד: cudaGetDeviceProperties ממלא מבנה cudaDeviceProp עם עשרות שדות שהדרייבר מדווח - כל מה שהתוכנית צריכה כדי להתאים את עצמה לחומרה. איך להכליל: אותה תוכנית רצה על כל כרטיס ומדפיסה את המספרים שלו; זו הדרך הנכונה לכתוב קוד נייד שמסתגל למספר ה-SM או ל-shared memory הזמין במקום להניח מספרים קשיחים.
פתרון תרגיל 4 - ה-kernel הראשון¶
// hello.cu
#include <cstdio>
__global__ void hello() {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
printf("Hello from thread %d (block %d, lane %d)\n",
tid, blockIdx.x, threadIdx.x);
}
int main() {
hello<<<2, 4>>>();
cudaDeviceSynchronize();
return 0;
}
$ nvcc -arch=sm_90 hello.cu -o hello
$ ./hello
Hello from thread 4 (block 1, lane 0)
Hello from thread 5 (block 1, lane 1)
Hello from thread 6 (block 1, lane 2)
Hello from thread 7 (block 1, lane 3)
Hello from thread 0 (block 0, lane 0)
Hello from thread 1 (block 0, lane 1)
Hello from thread 2 (block 0, lane 2)
Hello from thread 3 (block 0, lane 3)
בדיוק 8 שורות - thread אחד לכל אחד מ-2x4=8 הthreads. כשמריצים כמה פעמים, הסדר בין ה-blocks עשוי להשתנות (לפעמים block 0 קודם, לפעמים block 1) - ה-blocks רצים במקביל וללא סדר מובטח. בתוך block בודד ארבעת ה-threads שייכים לאותו warp ולכן נוטים לצאת ברצף, אבל אין להסתמך גם על זה.
כשמסירים את cudaDeviceSynchronize():
הפלט נעלם (או חלקי). ה-printf מdevice כותב לחוצץ שמתרוקן רק כשמסנכרנים או כשה-context נסגר מסודר; בלי הסנכרון, main חוזר וה-CPU יוצא לפני שה-GPU הספיק להריץ ולרוקן, והשורות אובדות.
למה זה עבד: <<<2, 4>>> השיק grid של 2 blocks עם 4 threads בכל אחד, וכל thread חישב אינדקס גלובלי ייחודי עם blockIdx.x * blockDim.x + threadIdx.x. איך להכליל: זו בדיוק נוסחת האינדוקס שבה נשתמש ב-0.3 כדי לתת לכל thread לעבד איבר אחר במערך; שם נחליף את ה-printf בכתיבה לזיכרון גלובלי דרך מצביע.
פתרון תרגיל 5 - תפיסת שגיאה עם CUDA_CHECK¶
$ nvcc -arch=sm_90 badcheck.cu -o badcheck
$ ./badcheck
CUDA error at badcheck.cu:14: invalid device ordinal
$ echo $?
1
התוכנית יצאה מיד עם קוד 1, והדפיסה את שם הקובץ, מספר השורה, וטקסט השגיאה - invalid device ordinal (התרגום של cudaErrorInvalidDevice).
למה זה עבד: כל קריאת Runtime API מחזירה cudaError_t; המקרו בדק אותו, ובמקרה של כישלון קרא ל-cudaGetErrorString והדפיס מיקום מדויק לפני exit. איך להכליל: עטפו כל קריאת Runtime API ב-CUDA_CHECK. עבור launches kernel, שאין להן ערך החזרה, בודקים בנפרד עם CUDA_CHECK(cudaGetLastError()) מיד אחרי הlaunch ו-CUDA_CHECK(cudaDeviceSynchronize()) כדי לתפוס שגיאות שקרו בזמן הריצה על הdevice - נעשה זאת החל מ-0.3.
פתרון תרגיל 6 (בונוס) - הכל בענן¶
ב-Colab, אחרי בחירת runtime מסוג GPU:
!nvidia-smi
...
| 0 Tesla T4 Off | 00000000:00:04.0 Off | 0 |
| N/A 38C P8 9W / 70W | 0MiB / 15360MiB | 0% Default |
בתא נפרד, כתיבת הקבצים והרצה:
%%writefile devinfo.cu
... (same code) ...
!nvcc -arch=sm_75 devinfo.cu -o devinfo && ./devinfo
Device 0: Tesla T4
Compute capability: 7.5
SM count: 40
Warp size: 32
Max threads / block: 1024
Shared mem / block: 48 KiB
Global (HBM) memory: 14.7 GiB
GPU boost clock: 1590 MHz
חישוב שיא ה-FP32 של T4 - שימו לב שדור Turing נושא 64 ליבות FP32 ל-SM, לא 128 כמו בדורות המודרניים:
T4: 40 SM x 64 FP32 cores/SM x 2 x 1.59 GHz = 8.1 TFLOPS
H100: 132 SM x 128 FP32 cores/SM x 2 x 1.98 GHz = 66.9 TFLOPS
ratio: 66.9 / 8.1 ~ 8.3x
ה-H100 חזק פי כ-8 מ-T4 ב-FP32 גולמי - וזה עוד לפני שמכניסים לתמונה את ה-Tensor Cores, שם הפער גדול בהרבה.
למה זה עבד: הקוד זהה לחלוטין; ההבדל היחיד הוא דגל ה--arch שהותאם ל-compute capability של T4 (7.5). איך להכליל: את כל הקורס אפשר להריץ בחינם ב-Colab בשינוי הזה בלבד. שימו לב שהנחת "128 ליבות FP32 ל-SM" נכונה ל-Ampere/Hopper אבל לא ל-Turing - תמיד ודאו את מספר הליבות לדור הספציפי לפני חישוב שיא, כי הוא משתנה בין ארכיטקטורות.