0.2 סביבת עבודה הרצאה

בשיעור 0.1 בנינו את המודל המנטלי של למה בכלל קיים GPU (מעבד גרפי) ומה ההבדל בין מעבד שמתוכנן ל-latency נמוך לבין מעבד שמתוכנן ל-throughput גבוה, וראינו את מודל ה-SIMT מלמעלה. עכשיו מגיע השיעור המעשי הראשון: נקים סביבת עבודה שרצה על כרטיס NVIDIA אמיתי, מקומית או בענן, ונוודא שהיא באמת עובדת. זה שיעור תשתית - לא נכתוב עדיין kernel רציני (זה מחכה לנו ב-0.3), אבל בלי הסביבה הזו אי אפשר להריץ שום דבר בהמשך הקורס. נלמד להבחין בין שני רכיבים שמתבלבלים בהם כל הזמן - הדרייבר של הכרטיס לעומת ה-CUDA Toolkit, נקרא כל שדה בפלט של nvidia-smi, נבין את הcompilation המפוצל של nvcc, ונכתוב שתי תוכניות קטנות: אחת ששואלת את הכרטיס מי הוא, ואחת שמדפיסה שלום מכל thread. לאורך כל הדרך ה-H100 (SXM) הוא הכרטיס לדוגמה שלנו.

שני חצאים של מחסנית התוכנה - הדרייבר מול ה-Toolkit

הטעות הכי נפוצה של מתחילים היא לחשוב ש"התקנתי CUDA" הוא דבר אחד. הוא לא. יש שני רכיבים נפרדים לגמרי, שמותקנים בנפרד, מתעדכנים בנפרד, ומגיעים מחבילות שונות:

• הדרייבר של NVIDIA - זה מה שמדבר עם החומרה. הוא כולל את מודול הליבה nvidia.ko (kernel module - מודול שרץ במצב ליבה, privileged mode, ומדבר ישירות עם הכרטיס), ואת הספרייה libcuda.so במרחב המשתמש שמממשת את ה-CUDA Driver API (הinterface ברמה הנמוכה, שכל הפונקציות שלו מתחילות ב-cu, למשל cuLaunchKernel). הדרייבר מותקן ברמת המערכת ומגיע מחבילת הדרייבר של NVIDIA.
• ה-CUDA Toolkit - זה מה שמפתחים איתו. הוא כולל את הcompiler nvcc, את הספרייה libcudart.so שמממשת את ה-CUDA Runtime API (הinterface ברמה הגבוהה, שכל הפונקציות שלו מתחילות ב-cuda, למשל cudaMalloc), את קובצי הכותרת (cuda_runtime.h), ואת ספריות הביצועים (cuBLAS, cuDNN). ה-Toolkit מגיע מחבילה נפרדת.

ההבחנה הזו איננה טכנית-פורמלית בלבד - היא הכלי האבחוני הכי חשוב שלכם. אם libcuda.so חסרה מהמערכת, זה כמעט תמיד אומר דבר אחד: הדרייבר לא מותקן (או מותקן שבור). לעומת זאת, אם nvcc חסר אבל nvidia-smi עובד, יש לכם דרייבר בלי Toolkit - מצב נפוץ מאוד, למשל על שרת ענן שבו מישהו התקין רק את הדרייבר. שימו לב לנקודה עדינה: libcuda.so מגיעה מהדרייבר, בעוד libcudart.so מגיעה מה-Toolkit. הן נשמעות כמעט זהות ומבלבלים ביניהן כל הזמן, אבל הן שני קבצים בשתי שכבות שונות.

     Our CUDA C++ code  (__global__, <<<>>>, cudaMalloc)
              |
              |  compiled with nvcc + gcc
              v
  +-------------------------------------------+
  |  CUDA Runtime API   (cuda*)   libcudart.so |   <-- from the Toolkit
  +-------------------------------------------+
              |
  +-------------------------------------------+
  |  CUDA Driver API    (cu*)     libcuda.so   |   <-- from the driver (user space)
  +-------------------------------------------+
              |   syscall / ioctl
  +-------------------------------------------+
  |  nvidia.ko    (kernel module, privileged)  |   <-- from the driver (kernel mode)
  +-------------------------------------------+
              |
           GPU hardware

הזרימה הזו היא כל המפה של פרק 5, ואנחנו נחזור אליה. לעת עתה מספיק לזכור: מלמעלה למטה - הקוד שלנו קורא ל-Runtime API, שיושב מעל ה-Driver API, שיושב מעל מודול הליבה, שנוגע בחומרה. אף בקשה מהhost לא מגיעה ל-GPU בלי לעבור את כל השרשרת הזו.

הבדיקה הראשונה - nvidia-smi

הפקודה הראשונה שמריצים על כל מכונה עם GPU היא nvidia-smi (ראשי תיבות של NVIDIA System Management Interface). זהו כלי שורת-פקודה שיושב מעל NVML (ספריית הניהול של NVIDIA, שממומשת ב-libnvml.so) ושואל את הכרטיס על מצבו. אם היא רצה - הדרייבר מותקן ועובד. הנה פלט טיפוסי מ-H100:

$ 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. |
|                                         |                        |               MIG M. |
|=========================================+========================+======================|
|   0  NVIDIA H100 80GB HBM3          On  |   00000000:53:00.0 Off |                    0 |
| N/A   32C    P0             71W / 700W  |       0MiB /  81559MiB  |      0%      Default |
|                                         |                        |             Disabled |
+-----------------------------------------+------------------------+----------------------+

בואו נקרא כל שדה, כי כל אחד מהם יחזור בהמשך הקורס:

• שורת הכותרת מפרידה בין שני מספרי גרסה שמבלבלים ביניהם: Driver Version הוא גרסת הדרייבר (כאן 550.54.15), ו-CUDA Version הוא גרסת ה-CUDA המקסימלית שהדרייבר הזה תומך בה (12.4). שימו לב שזו לא בהכרח גרסת ה-Toolkit המותקנת - זו התקרה שהדרייבר מאפשר. אפשר להתקין Toolkit ישן יותר, אבל לא חדש יותר מהמספר הזה.
• GPU - האינדקס של הכרטיס (0). במכונה עם שמונה כרטיסים תראו 0 עד 7.
• Name - דגם הכרטיס, כאן NVIDIA H100 80GB HBM3.
• Persistence-M - מצב התמדה (persistence mode). כשהוא On, הדרייבר נשאר טעון גם כשאין תהליך שמשתמש ב-GPU, מה שחוסך כמה שניות של אתחול בכל הרצה. במחשב אישי הוא בדרך כלל Off; בשרתים מדליקים אותו עם nvidia-smi -pm 1.
• Bus-Id - כתובת ה-PCIe של הכרטיס (00000000:53:00.0). שימושי כדי לזהות כרטיס פיזי ספציפי במכונה עם הרבה כרטיסים.
• Disp.A - האם מחובר צג לכרטיס (Display Active). על כרטיס מרכז-נתונים כמו H100 זה תמיד Off.
• Volatile Uncorr. ECC - מונה שגיאות זיכרון שלא תוקנו מאז האתחול. 0 זה טוב; מספר שגדל הוא סימן לחומרה תקולה.
• Fan - מהירות מאוורר. על H100 (כרטיס פסיבי בשרת מקורר) זה N/A - אין לו מאוורר משלו.
• Temp - טמפרטורת ה-GPU במעלות צלזיוס (32C - כרטיס בטל). תחת עומס תראו 60-80 מעלות; מעל סף מסוים הכרטיס יעשה throttling ויוריד תדר.
• Perf - מצב הביצועים (performance state), מ-P0 (הכי פעיל, תדר מלא) ועד P12 (הכי חסכוני). כרטיס בטל שממתין לעבודה יכול להיות ב-P0 או במצב חסכוני יותר.
• Pwr:Usage/Cap - צריכת ההספק הרגעית מול תקרת ההספק. כאן 71W / 700W: הכרטיס צורך 71 ואט בבטלה, והתקרה שלו היא 700 ואט (ה-TDP של H100 SXM). את התקרה אפשר לשנות עם -pl.
• Memory-Usage - זיכרון ה-HBM שבשימוש מול הסך הכל. כאן 0MiB / 81559MiB - כלומר כ-80GB פנויים. שימו לב שהמספר 81559MiB שנציג פה שונה מעט מהזיכרון ש-cudaGetDeviceProperties מדווח בתוך תוכנית, כי חלק שמור לדרייבר.
• GPU-Util - אחוז הזמן שבו ה-GPU היה עסוק בהרצת kernel בחלון הדגימה האחרון. זהו מדד גס מאוד - 100% אומר שרץ kernel כלשהו, לא שכל יחידות הביצוע נוצלו. ניצול אמיתי נמדוד עם Nsight בפרק 9.
• Compute M. - מצב החישוב (compute mode). Default אומר שכמה תהליכים יכולים לחלוק את הכרטיס; Exclusive_Process נועל את הכרטיס לתהליך יחיד.
• MIG M. - האם Multi-Instance GPU פעיל. MIG מחלק כרטיס אחד לכמה כרטיסים וירטואליים מבודדים. כאן Disabled.

לצלול פנימה - nvidia-smi -q ורשימת התהליכים

הטבלה שראינו היא תקציר. לפירוט מלא, שכולל הרבה יותר שדות, מריצים את nvidia-smi -q (query). הפלט ארוך ומחולק לקטעים; הנה קטע מייצג:

$ nvidia-smi -q

==============NVSMI LOG==============
Product Name                          : NVIDIA H100 80GB HBM3
GPU UUID                              : GPU-4a1b2c3d-...
    Compute Mode                      : Default
    FB Memory Usage
        Total                         : 81559 MiB
        Used                          : 0 MiB
        Free                          : 81559 MiB
    Temperature
        GPU Current Temp              : 32 C
        GPU Slowdown Temp             : 85 C
        GPU Shutdown Temp             : 90 C
    Power Readings
        Power Draw                    : 71.05 W
        Current Power Limit           : 700.00 W
    Clocks
        SM                            : 345 MHz
        Memory                        : 2619 MHz

כאן רואים דברים שהטבלה הקצרה הסתירה: את ה-UUID (מזהה ייחודי של הכרטיס, יציב יותר מהאינדקס), את ספי ההאטה והכיבוי התרמיים, ואת התדרים הנוכחיים. שימו לב שתדר ה-SM בבטלה (345 MHz) נמוך בהרבה מתדר ה-boost המלא (סביב 1980 MHz) - הכרטיס מוריד תדר כשאין עבודה.

כדי לראות מי משתמש בכרטיס כרגע, מריצים nvidia-smi ומסתכלים בטבלת התהליכים בתחתית, או nvidia-smi pmon לניטור מתמשך:

+-----------------------------------------------------------------------------------------+
| Processes:                                                                              |
|  GPU   GI   CI        PID   Type   Process name                              GPU Memory |
|        ID   ID                                                               Usage      |
|=========================================================================================|
|  No running processes found                                                             |
+-----------------------------------------------------------------------------------------+

כשתריצו תוכנית CUDA, ה-PID שלה יופיע כאן עם כמות הזיכרון שהיא הקצתה על הכרטיס. זו הדרך הכי מהירה לתפוס תהליך "זומבי" ששכח לשחרר זיכרון ותופס לכם את כל ה-HBM.

אזהרה חשובה לכל מי שכותב כלים: פורמט הטקסט של nvidia-smi אינו יציב - NVIDIA לא מבטיחה שהוא לא ישתנה בין גרסאות. אל תכתבו סקריפט שמנתח (parse) את הפלט הטקסטואלי הזה. אם אתם צריכים לקרוא טלמטריה בתוך כלי, פנו ישירות ל-NVML C API (או לעטיפות שלו כמו pynvml בפייתון), שהוא הinterface היציב. לפלט מובנה יותר אפשר גם nvidia-smi -q -x שמוציא XML.

הcompiler - nvcc וה-compilation המפוצל

עכשיו לצד ה-Toolkit. הפקודה שמאמתת אותו היא:

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Tue_Feb_27_16:19:38_PST_2024
Cuda compilation tools, release 12.4, V12.4.99
Build cuda_12.4.r12.4/compiler.34097967_0

השורה החשובה היא release 12.4 - זו גרסת ה-Toolkit המותקנת. אם הפקודה נכשלת ב-command not found אבל nvidia-smi עבד, יש לכם דרייבר בלי Toolkit, וצריך להתקין את ה-Toolkit או להוסיף את /usr/local/cuda/bin ל-PATH.

השם nvcc מטעה: זה לא compiler מונוליטי אלא compiler-מנהל (compiler driver) - מתזמר שמפעיל כמה כלים מתחתיו. ההבחנה הזו קריטית, וה-Glossary מדגיש אותה: compiler-מנהל כמו nvcc או gcc איננו מנהל-device חומרתי כמו הדרייבר של NVIDIA. אותה מילה "driver", שני דברים שונים לגמרי.

מה nvcc עושה בפועל הוא compilation מפוצל (split compilation). קובץ .cu מכיל שני סוגי קוד מעורבבים: קוד host שרץ על ה-CPU, וקוד device שרץ על ה-GPU - הפונקציות המסומנות __global__ ו-__device__. nvcc מפריד ביניהם:

        program.cu
            |
     +------+------------------------+
     |                               |
  host code                     device code
     |                               |
   gcc                            nvcc  ->  PTX  ->  SASS
     |                               |
     +---------------+---------------+
                     |
              single executable "fat binary"
        (host ABI + PTX and/or SASS embedded inside)

את קוד הhost nvcc מוסר לcompiler המערכת - בדרך כלל gcc - בדיוק כמו כל תוכנית C++. את קוד הdevice הוא מהדר בעצמו: קודם ל-PTX (שפת אסמבלי וירטואלית, ניידת בין דורות), ואז ל-SASS (האסמבלי האמיתי של הארכיטקטורה הספציפית). התוצאה היא קובץ הרצה אחד, "fat binary" (fatbin), שהוא בינארי host רגיל שגם נושא בתוכו את קוד ה-GPU. בגלל שהוא בינארי host רגיל אפשר לפרק אותו עם readelf הרגיל, או עם כלי הבינארי של CUDA (cuobjdump, nvdisasm) שנכיר בפרק 4. את כל שרשרת הcompilation הזו נפרק לגמרי בפרק 4 - כאן מספיק לדעת שכשאנחנו מריצים nvcc file.cu -o file, שני compilers עובדים מאחורי הקלעים.

הדגל שקובע לאיזו ארכיטקטורה לקמפל הוא -arch. ל-H100 נכתוב -arch=sm_90 (ה-compute capability של Hopper). הערך sm_90 קובע איזה SASS ייווצר; ערך compute_90 היה קובע איזו גרסת PTX תוטמע. אם לא נציין -arch, nvcc יבחר ברירת מחדל ישנה, ולפעמים שמרנית מדי - בקורס הזה נציין ארכיטקטורה תמיד.

למי שאין GPU - Colab וענן

לא לכולם יש כרטיס NVIDIA במחשב, וזה בסדר גמור - כל הקורס ניתן לתרגול בענן. הנה הדרכים המהירות לקבל shell עם GPU:

• Google Colab - הדרך החינמית והמהירה ביותר. נכנסים ל-colab.research.google.com, פותחים notebook, ובתפריט Runtime -> Change runtime type בוחרים GPU. מקבלים בדרך כלל כרטיס T4 (מדור Turing, compute capability 7.5, 16GB) עם CUDA כבר מותקן. אפשר להריץ nvidia-smi בתא עם סימן קריאה לפניו (!nvidia-smi), ולקמפל קובץ .cu ישירות עם !nvcc. ה-T4 חלש בהרבה מ-H100, אבל מצוין ללמידה - כל הקוד בקורס ירוץ עליו (עם -arch=sm_75 במקום sm_90).
• Modal - פלטפורמת ענן שמריצה פונקציות פייתון על GPU לפי דרישה, כולל H100 ו-A100. משלמים לפי שנייה. מגדירים את הכרטיס בקוד (gpu="H100") ומריצים. מתאים למי שרוצה בדיוק את הכרטיס של הקורס.
• Lambda ו-מכונות GPU בענן (AWS, GCP, Azure) - שוכרים מכונה וירטואלית עם GPU לפי שעה, מתחברים ב-SSH, ומקבלים shell מלא. גמיש ביותר אבל דורש התקנת דרייבר ו-Toolkit (או שימוש ב-image מוכן עם CUDA).

לכל אורך הקורס נציין את המספרים של H100, אבל בכל תרגיל תוכלו להריץ את אותו הקוד על T4 בחינם ולראות את המספרים של הכרטיס שלכם. ההבדל היחיד בפועל הוא דגל ה--arch.

לשאול את הכרטיס מי הוא - cudaGetDeviceProperties

nvidia-smi נותן תמונה חיצונית. כדי לשאול את הכרטיס על המאפיינים הארכיטקטוניים שלו מתוך קוד, קוראים ל-cudaGetDeviceProperties. זו הגרסה המינימלית של הכלי deviceQuery שמגיע עם דוגמאות ה-CUDA. נכתוב אותה בעצמנו, ותוך כדי נכיר את מקרו בדיקת השגיאות CUDA_CHECK - כלי שנשתמש בו בכל הקורס. כל קריאה ל-Runtime API מחזירה cudaError_t, ומי שלא בודק אותה מגלה באגים מאוחר מדי. המקרו עוטף כל קריאה, ואם היא נכשלה - מדפיס את שם השגיאה, את הקובץ ואת השורה, ויוצא.

// 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("  Max threads / SM:     %d\n", p.maxThreadsPerMultiProcessor);
        printf("  Shared mem / block:   %zu KiB\n", p.sharedMemPerBlock / 1024);
        printf("  Shared mem / SM:      %zu KiB\n", p.sharedMemPerMultiprocessor / 1024);
        printf("  Registers / block:    %d\n", p.regsPerBlock);
        printf("  Global (HBM) memory:  %.1f GiB\n",
               p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
        printf("  Memory bus width:     %d bit\n", p.memoryBusWidth);
        printf("  L2 cache:             %d MiB\n", p.l2CacheSize / (1024 * 1024));
        printf("  GPU boost clock:      %.0f MHz\n", p.clockRate / 1000.0);
    }
    return 0;
}

מהדרים ומריצים:

$ nvcc -arch=sm_90 devinfo.cu -o devinfo
$ ./devinfo

והנה הפלט על H100:

CUDA devices found: 1

Device 0: NVIDIA H100 80GB HBM3
  Compute capability:   9.0
  SM count:             132
  Warp size:            32
  Max threads / block:  1024
  Max threads / SM:     2048
  Shared mem / block:   48 KiB
  Shared mem / SM:      228 KiB
  Registers / block:    65536
  Global (HBM) memory:  79.1 GiB
  Memory bus width:     5120 bit
  L2 cache:             50 MiB
  GPU boost clock:      1980 MHz

כל מספר כאן הוא עובדה שנחזור אליה בקורס. ה-H100 הוא 132 SM (מעבד רב-זרמי - זו יחידת המקבול האמיתית, כפי שראינו ב-0.1), warp של 32 threads, עד 1024 threads ל-block ועד 2048 threads (64 warps) שיושבים בו-זמנית על SM יחיד. שימו לב ל-Shared mem / block שהוא 48 KiB בברירת מחדל, בעוד ה-SM עצמו נושא 228 KiB - הפער הזה (opt-in לזיכרון גדול יותר) יהיה נושא שלם בפרק 3. שימו לב גם שהזיכרון המדווח כאן, 79.1 GiB, שונה מעט מ-81559MiB של nvidia-smi - השני מודד את הזיכרון הפיזי הכולל, הראשון את מה שזמין לתוכניות אחרי ששמור לדרייבר.

בואו נשתמש במספרים כדי לחשב את שיא ה-FP32 של הכרטיס, בדיוק כמו שנעשה בפרק 7. ל-H100 יש 128 ליבות FP32 (CUDA Cores) בכל SM, וכל ליבה מבצעת פעולת FMA (כפל-חיבור, שתי פעולות נקודה-צפה) בכל מחזור שעון:

132 SM x 128 FP32 cores/SM x 2 FLOP/FMA x 1.98 GHz
   = 16,896 cores x 2 x 1.98e9
   = 66.9 TFLOPS  (FP32, without Tensor Cores)

66.9 TFLOPS - בדיוק המספר שה-Glossary נותן ל-H100. שווה לעצור על זה: את הנתונים לחישוב הזה קיבלנו כולם מ-cudaGetDeviceProperties ומידע ארכיטקטוני, בלי להסתכל בשום דף מוצר. זו המיומנות שהקורס בונה.

מבחן העשן - שלום מכל thread

הבדיקה האחרונה: לקמפל ולהריץ kernel אמיתי, כדי לוודא שכל שרשרת הcompilation וההרצה עובדת מקצה לקצה. זה עדיין לא kernel מעניין - הוא רק מדפיס. אבל הוא מוכיח שהקוד שלנו באמת רץ על ה-GPU ולא על ה-CPU.

// 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>>>();          // 2 blocks, 4 threads each = 8 threads
    cudaDeviceSynchronize();    // waits for the GPU to finish before the host exits
    return 0;
}

שימו לב לשלושה דברים שכבר ראינו במפה של 0.1, ושנרחיב בהם ב-0.3. הפונקציה מסומנת __global__ - זהו kernel, פונקציה שהhost משיק וdevice מריץ. מחזירה void תמיד: kernel לא "מחזיר" ערך, הוא כותב תוצאות דרך מצביעים (כאן הוא רק מדפיס). והlaunch נעשית עם התחביר <<<2, 4>>> - שני מספרים שקובעים 2 blocks של 4 threads. הביטוי blockIdx.x * blockDim.x + threadIdx.x הוא נוסחת האינדוקס הקנונית שכל thread מחשב בה מי הוא; נשבור אותה לגמרי ב-0.3.

מהדרים ומריצים:

$ 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)

שמונה שורות, אחת מכל thread - הצלחנו. שימו לב שהסדר לא בהכרח 0 עד 7. ה-blocks רצים במקביל וללא סדר מובטח, ולכן ייתכן ש-block 1 ידפיס לפני block 0. זו לא תקלה - זו בדיוק המקביליות של ה-GPU, ולעולם אל תסתמכו על סדר בין threads או בין blocks. אם הפלט ריק לגמרי, כמעט תמיד הסיבה היא ש-cudaDeviceSynchronize() חסר: הhost יצא לפני שה-GPU הספיק לרוקן את חוצץ ה-printf. אם הקימפול נכשל על sm_90, ה-Toolkit שלכם ישן מדי ל-Hopper (צריך CUDA 12.0 ומעלה) - על Colab עם T4 השתמשו ב--arch=sm_75.

סיכום

• מחסנית ה-GPU מתחלקת לשני רכיבים נפרדים: הדרייבר (מספק את nvidia.ko במצב ליבה ואת libcuda.so במרחב המשתמש) וה-CUDA Toolkit (מספק את nvcc, את libcudart.so ואת הספריות); הם מותקנים ומתעדכנים בנפרד.
• היעדר libcuda.so פירושו כמעט תמיד שהדרייבר לא מותקן; היעדר nvcc בזמן ש-nvidia-smi עובד פירושו דרייבר בלי Toolkit.
• הפקודה nvidia-smi מאמתת שהדרייבר עובד ומציגה טבלה שבה כל שדה משמעותי: אינדקס, דגם, persistence, Bus-Id, טמפרטורה, מצב ביצועים, הספק מול תקרה, שימוש בזיכרון, ניצול, מצב חישוב ו-MIG; nvidia-smi -q נותן פירוט מלא ורשימת תהליכים.
• פורמט הטקסט של nvidia-smi אינו יציב - כלים אמינים צריכים לפנות ל-NVML C API (או ל-pynvml), לא לנתח את הפלט הטקסטואלי.
• הפקודה nvcc --version מאמתת את ה-Toolkit; nvcc הוא compiler-מנהל שמבצע compilation מפוצל - gcc על קוד הhost ו-nvcc על קוד הdevice (PTX ואז SASS), ומייצר "fat binary" יחיד.
• מי שאין לו כרטיס NVIDIA יכול לתרגל את כל הקורס ב-Google Colab בחינם (כרטיס T4, -arch=sm_75), ב-Modal (H100/A100 לפי שנייה), או במכונות GPU בענן.
• הפונקציה cudaGetDeviceProperties שואלת את הכרטיס על מאפייניו - מספר ה-SM, ה-compute capability, warp size, threads מקסימלי ל-block, shared memory וגודל ה-HBM; מהמספרים האלה חישבנו את שיא ה-FP32 של H100 (66.9 TFLOPS).
• מקרו CUDA_CHECK שעוטף כל קריאת Runtime API הוא הרגל חובה - הוא הופך שגיאה שקטה להודעה עם קובץ ושורה.
• מבחן העשן - kernel __global__ פשוט שמושק ב-<<<2,4>>> ומודפס עם cudaDeviceSynchronize() - מוכיח שכל השרשרת עובדת; הסדר בין threads אינו מובטח, וזו תכונה של המקביליות ולא באג.