5.2 הדרייבר nvidia.ko ו libcuda הרצאה

בשיעור 5.1 הכרנו את הדרייבר של NVIDIA (מנהל הdevice / driver) כשכבת התיווך שדרכה, ורק דרכה, קוד על ה-host (host / ה-CPU) מגיע אל ה-GPU (מעבד גרפי / כרטיס מסך). ראינו שהדרייבר אינו מונוליט אחד אלא נחתך לרוחב גבול ה-user/kernel לשני חצאים. השיעור הזה פותח את שני החצאים ומסתכל להם פנימה: מצד אחד nvidia.ko - מודול ליבה (kernel module) מיוחס שרץ בתוך ליבת לינוקס ונוגע פיזית בסיליקון, ומצד שני libcuda.so - ספרייה סגורה במרחב המשתמש שמממשת את ה-Driver API. נעקוב אחרי בקשה בודדת מקוד המשתמש עד לחומרה ובחזרה, נראה בדיוק היכן היא חוצה את גבול ה-kernel דרך ioctl, נכיר את רכיב ה-UVM (Unified Virtual Memory - זיכרון וירטואלי מאוחד) ואת ה-JIT של הדרייבר שמהדר PTX ל-SASS בזמן טעינה, ונחתום בהבחנה שמבלבלת חצי מהעולם: ההבדל בין הדרייבר לבין ה-Toolkit. כאן נשען ישירות על מה שלמדתם על ליבת מערכת ההפעלה בקורס ליבת-המחשב - device drivers, מרחב ליבה מול מרחב משתמש, ו-syscalls.

הדרייבר כמתווך יחיד - the sole mediator

הרעיון המרכזי, שחוזר על עצמו לאורך כל הפרק, הוא שהדרייבר הוא נקודת המעבר היחידה אל ה-GPU. אף שורת קוד ביישום שלכם אינה כותבת לרגיסטר של הכרטיס, אינה מתחילה העברת DMA ואינה משגרת grid (רשת) של blocks (בלוקים) ישירות. כל פעולה כזו נארזת כבקשה, עוברת מטה דרך שכבות התוכנה, וחוצה בסופו של דבר את גבול ה-kernel אל תוך nvidia.ko, שהוא הרכיב היחיד המורשה לגעת בחומרה. זהו בדיוק אותו עיקרון שהכרתם בליבת-המחשב: תהליך במרחב המשתמש אינו ניגש לdevice ישירות אלא מבקש מהליבה, והליבה - דרך ה-device driver המתאים - מבצעת בפועל.

הדרייבר נחתך לשניים סביב גבול ה-user/kernel:

  User space (user space, unprivileged)      Kernel space (kernel space, privileged)
  ------------------------------------        ---------------------------------
  libcudart.so   (Runtime API)                nvidia.ko          (core)
  libcuda.so     (Driver API)   <===ioctl===> nvidia-uvm.ko      (UVM)
                                              nvidia-modeset.ko  (display)
                                              nvidia-drm.ko      (DRM)
  --- where your code runs ---                --- where they touch the hardware ---

שימו לב לחלוקת התפקידים. libcuda.so יושבת במרחב המשתמש, לא מיוחסת, ותפקידה לתרגם קריאות Driver API לבקשות ולהעביר אותן מטה. nvidia.ko יושבת במרחב הליבה, מיוחסת, ותפקידה לקבל את הבקשות ולבצע אותן מול הסיליקון. שני החצאים יחד הם "הדרייבר", אבל הם קבצים נפרדים, במרחבי הרצה נפרדים, עם רמות הרשאה נפרדות. הבחנה זו היא לב השיעור.

מודול הליבה - nvidia.ko

nvidia.ko הוא, כלשון ה-glossary, "קובץ מודול ליבה בינארי בליבם של דרייברי ה-GPU של NVIDIA ללינוקס". נפרק את המשפט הזה.

הסיומת .ko היא ראשי תיבות של kernel object - הפורמט של מודול ליבה נטען (loadable kernel module) בלינוקס. זהו קובץ ELF, בן דוד של קובץ .o רגיל, אבל במקום להתקשר לתוך תוכנית משתמש הוא נטען לתוך ליבת מערכת ההפעלה הרצה ומריץ את הקוד שלו במרחב הליבה. מרגע שנטען, הקוד שלו רץ במצב מיוחס (ring 0 / privileged mode) - אותה רמת הרשאה של הליבה עצמה. זו הסיבה שהוא, ורק הוא, יכול לתכנת את בקר ה-DMA של הכרטיס, למפות את רגיסטרי ה-MMIO שלו, ולטפל בפסיקות (interrupts) שה-GPU שולח.

בפועל אין מודול אחד אלא משפחה של מודולים שנטענים יחד:

מודול	תפקיד
nvidia.ko	הליבה: ניהול הdevice, זיכרון, dispatch עבודה לחומרה
nvidia-uvm.ko	רכיב ה-Unified Virtual Memory (נעמיק בהמשך)
nvidia-modeset.ko	ניהול מצבי תצוגה (רלוונטי לכרטיסים גרפיים, פחות ל-H100 בשרת)
nvidia-drm.ko	אינטגרציה עם תשתית ה-DRM (Direct Rendering Manager) של הליבה

טעינה - modprobe. מודול נטען לליבה בעזרת modprobe, שגם פותר תלויות בין מודולים אוטומטית (בניגוד ל-insmod הגולמי). בדרך כלל הטעינה קורית אוטומטית באתחול, אבל אפשר לעשות זאת ידנית:

sudo modprobe nvidia          # resolves dependencies and loads nvidia.ko and what depends on it

בדיקה - lsmod. כדי לראות אילו מודולים טעונים כרגע מריצים lsmod (שהוא בעצם קריאה יפה של /proc/modules):

lsmod | grep nvidia

nvidia_uvm           1560576  0
nvidia_drm             77824  4
nvidia_modeset       1355776  8 nvidia_drm
nvidia               54587392 322 nvidia_uvm,nvidia_modeset

קִראו את הטבלה: העמודה הראשונה היא שם המודול, השנייה גודלו בבתים, השלישית מונה השימוש (כמה משתמשים בו כרגע), והרביעית - מי תלוי בו. שימו לב ש-nvidia (השורה התחתונה) הוא הבסיס, ו-nvidia_uvm ו-nvidia_modeset תלויים בו. מונה השימוש הגבוה (322) על המודול הראשי משקף context-ים, mappings ותהליכים פעילים שמחזיקים בו כרגע.

מטא-נתונים - modinfo. הפקודה modinfo שולפת את מחרוזות המטא-נתונים המוטמעות בקובץ ה-.ko, כולל הגרסה והרישיון:

modinfo nvidia | head -n 5

filename:       /lib/modules/6.8.0-40-generic/kernel/drivers/video/nvidia.ko
version:        550.54.15
supported:      external
license:        Dual MIT/GPL

הרישיון Dual MIT/GPL הוא סימן ההיכר של גרסת הקוד הפתוח. בגרסה הקניינית הישנה שדה הרישיון היה NVIDIA, וטעינת המודול "צבעה" את הליבה כ-tainted (מזוהמת) בגלל קוד לא-פתוח. עוד על כך בסעיף הבא.

הודעות - dmesg. מודול ליבה אינו יכול להדפיס ל-stdout; הוא כותב ל-kernel ring buffer, שנקרא עם dmesg. כאן מופיעות הודעות האתחול של הדרייבר, שגיאות XID (שגיאות חומרה של ה-GPU), ואירועי טעינה. ההודעות של NVIDIA מסומנות בקידומת NVRM (NVIDIA Resource Manager):

sudo dmesg | grep -iE 'nvidia|nvrm'

[    5.234567] nvidia: loading out-of-tree module taints kernel.
[    5.245678] NVRM: loading NVIDIA UNIX Open Kernel Module for x86_64  550.54.15  Release Build
[    5.356789] nvidia 0000:17:00.0: enabling device (0000 -> 0002)
[    5.467890] nvidia-uvm: Loaded the UVM driver, major device number 235.

השורה השנייה מאשרת שזו הגרסה הפתוחה (Open Kernel Module), והאחרונה מראה ש-nvidia-uvm רשם לעצמו major device number 235 - נחזור לזה כשנדבר על קובצי הdevice.

קוד פתוח - open-gpu-kernel-modules. נכון להיום, החצי של הליבה שוחרר כקוד פתוח תחת הפרויקט NVIDIA/open-gpu-kernel-modules בגיטהאב (רישיון Dual MIT/GPL). זו נקודת מפנה: היסטורית כל הדרייבר היה סגור. עבור דורות Turing ומעלה - כלומר בוודאי עבור ה-H100 (Hopper) שלנו - המודול הפתוח הוא כיום ברירת המחדל המומלצת. חשוב להדגיש: רק חצי ה-kernel נפתח. החצי במרחב המשתמש, ובראשו libcuda.so, נשאר סגור לחלוטין. זו לא סתירה - הרבה מהלוגיקה "החכמה" של הדרייבר חיה דווקא במרחב המשתמש.

הספרייה במרחב המשתמש - libcuda.so

מהעבר השני של גבול ה-kernel יושבת libcuda.so - "קובץ ה-shared object הבינארי שמממש את ה-CUDA Driver API בלינוקס". זהו החצי הבלתי-מיוחס של הדרייבר.

הפונקציות שהיא חושפת נושאות את הקידומת cu (למשל cuInit, cuMemAlloc, cuModuleLoad, cuLaunchKernel) - להבדיל מקידומת cuda של ה-Runtime API. שני מאפיינים חשובים:

• קישור דינמי. לעולם לא מקשרים את libcuda.so סטטית; היא תמיד מקושרת דינמית ונטענת בזמן ריצה. הסיבה עמוקה: הספרייה חייבת להתאים לכרטיס ולגרסת הדרייבר של המכונה שעליה רצים בפועל, ואלה ידועים רק בזמן ריצה, לא בזמן compilation.
• מגיעה עם הדרייבר, לא עם ה-Toolkit. libcuda.so היא חלק מחבילת הדרייבר של NVIDIA, ולא מ-CUDA Toolkit. מכאן אבחון מעשי חשוב: אם libcuda.so חסרה מהמערכת, המשמעות כמעט תמיד היא שהדרייבר לא הותקן כראוי.

נבדוק את נוכחות שתי הספריות ואת מיקומן:

ldconfig -p | grep -E 'libcuda|libcudart'

    libcudart.so.12 (libc6,x86-64) => /usr/local/cuda/lib64/libcudart.so.12
    libcuda.so.1 (libc6,x86-64) => /lib/x86_64-linux-gnu/libcuda.so.1
    libcuda.so (libc6,x86-64) => /lib/x86_64-linux-gnu/libcuda.so

הבחינו בפרט עדין ומאיר עיניים: ה-soname של הספרייה הוא libcuda.so.1 - המספר הראשי נשאר 1 וקבוע על פני כל גרסאות הדרייבר. זו הבטחת ה-ABI היציב שמאפשרת תאימות קדימה: יישום שקושר מול libcuda.so.1 ימשיך לרוץ גם כשמתקינים דרייבר חדש יותר, כי הקובץ הפיזי (למשל libcuda.so.550.54.15) מתחלף אבל ה-soname לא. לעומת זאת ה-soname של ספריית ה-Runtime הוא libcudart.so.12, ומספרו עוקב אחרי הגרסה הראשית של ה-Toolkit.

סגורה, ויש חלופות. מימוש ה-Driver API סגור. עם זאת קיימים ניסיונות קוד-פתוח למֵמֵש אותו מבחוץ - LibreCuda ו-tinygrad הם השניים הבולטים - בדיוק משום שהחצי הזה של הדרייבר לא נפתח יחד עם מודול הליבה.

נתיב הבקשה - the request path via ioctl

עכשיו נחבר את שני החצאים. נעקוב אחרי cudaMalloc בודד, מקוד המשתמש עד לזיכרון ה-GPU, ונראה בדיוק היכן הבקשה חוצה אל הליבה. הנתיב הוא:

  User code
     |  cudaMalloc(&p, n)
     v
  libcudart.so   (Runtime API, user mode)   -- converts to a Driver API call
     |  cuMemAlloc(...)
     v
  libcuda.so     (Driver API,  user mode)   -- packs a request to nvidia.ko
     |  ioctl(fd, NV_ESC_..., &params)
  === user / kernel boundary (syscall trap) ===
     |
     v
  nvidia.ko      (core, kernel mode)        -- executes against the hardware
     |  MMIO / DMA / interrupts
     v
  GPU hardware (device):  80 GiB HBM3 of H100

הצומת המעניין הוא המעבר מ-libcuda.so אל nvidia.ko. הוא אינו קריאת פונקציה רגילה - הוא חוצה גבול הרשאות, ולכן חייב לעבור דרך syscall. המנגנון שבו NVIDIA משתמשת הוא ioctl (input/output control) - אותה קריאת מערכת גנרית לשליטה בהתקנים שהכרתם בליבת-המחשב. הרעיון: הדרייבר חושף קובצי device (device nodes) תחת /dev, קוד המשתמש פותח אותם ומנפיק אליהם ioctl עם קוד פקודה ומבנה פרמטרים, והליבה מפענחת ומבצעת.

קובצי הdevice שהדרייבר יוצר:

קובץ device	תפקיד
/dev/nvidiactl	ערוץ בקרה גלובלי (ניהול context, הקצאות)
/dev/nvidia0	ה-GPU הפיזי הראשון (nvidia1 לשני, וכן הלאה)
/dev/nvidia-uvm	ערוץ ה-Unified Virtual Memory
/dev/nvidia-uvm-tools	ערוץ כלים/פרופיילינג של UVM

אפשר לראות את הצומת הזה בעיניים ממש, בעזרת strace, שמיירט את קריאות המערכת של תוכנית. נריץ תוכנית CUDA פשוטה תחת strace ונסנן ל-ioctl ו-openat:

strace -f -e trace=openat,ioctl ./vecadd 2>&1 | grep -E 'nvidia|ioctl' | head

openat(AT_FDCWD, "/dev/nvidiactl", O_RDWR)      = 3
ioctl(3, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2a, 0x20), 0x7ffd...) = 0
openat(AT_FDCWD, "/dev/nvidia0", O_RDWR)        = 4
ioctl(4, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2b, 0x38), 0x7ffd...) = 0
openat(AT_FDCWD, "/dev/nvidia-uvm", O_RDWR)     = 5
ioctl(5, _IOC(_IOC_READ|_IOC_WRITE, 0x00, 0x02, 0x110), 0x7ffd...) = 0

הביטו על ה-magic number של ה-ioctl על שני הקבצים הראשונים: 0x46, שהוא קוד ה-ASCII של האות F. זהו קבוע ה-NV_IOCTL_MAGIC שהדרייבר בחר. כל שיחה בין libcuda.so ל-nvidia.ko היא, בסופו של דבר, זרם של קריאות ioctl כאלה. זהו הpipe הצר שדרכו עוברת כל אינטראקציה עם ה-GPU - הקצאת זיכרון, העתקת נתונים, dispatch kernel (קרנל - פונקציית ה-GPU) והמתנה לסיומו.

הערה שמונעת בלבול: cudaMemcpy של מערך גדול אינו מעביר את הבתים בתוך ה-ioctl עצמו. ה-ioctl רק מתכנת את מנוע ה-DMA של הכרטיס, וזה מושך את הנתונים ישירות מזיכרון ה-host דרך ה-PCIe/NVLink. ה-ioctl הוא ערוץ הבקרה; הנתונים זורמים בערוץ נפרד ומהיר בהרבה.

רכיב הזיכרון הווירטואלי המאוחד - UVM

שמתם לב שב-strace נפתח גם /dev/nvidia-uvm נפרד. זהו רכיב מיוחד של הדרייבר, ממומש במודול הליבה nvidia-uvm.ko, שאחראי על Unified Virtual Memory - זיכרון וירטואלי מאוחד. הרעיון: מרחב כתובות וירטואלי אחד המשותף ל-host ול-device, כך שאותו מצביע תקף בשני הצדדים, והדרייבר מהגר דפים (pages) בין זיכרון ה-CPU ל-HBM של ה-GPU לפי הצורך.

זהו המנגנון שמאחורי cudaMallocManaged:

float* data;
CUDA_CHECK(cudaMallocManaged(&data, n * sizeof(float)));  // a single pointer for both sides
for (int i = 0; i < n; ++i) data[i] = 1.0f;               // the CPU writes
vecAdd<<<blocks, threads>>>(data, data, data, n);          // the GPU reads/writes - same pointer
CUDA_CHECK(cudaDeviceSynchronize());
printf("%f\n", data[0]);                                    // the CPU reads again
CUDA_CHECK(cudaFree(data));

כאן data הוא מצביע יחיד שה-CPU וה-GPU משתמשים בו בלי cudaMemcpy מפורש. איך זה עובד מתחת למכסה: כשה-GPU ניגש לדף שנמצא כרגע בזיכרון ה-host, החומרה יורה page fault, המטופל על ידי nvidia-uvm.ko. המודול משהה את ה-warp (קבוצת 32 threads) הרלוונטי, מהגר את הדף אל ה-HBM של ה-GPU, מעדכן את טבלאות הדפים, וממשיך. זו בדיוק אותה מכניקה של demand paging מ-CPU שהכרתם בליבת-המחשב, אלא שכאן ה"פסיקה" מגיעה מה-GPU דרך המודול הזה. הנוחות עולה בlatency (latency): page fault הוא יקר, ולכן בקוד ביצועים גבוה נעדיף לרוב הקצאה מפורשת עם cudaMalloc והעתקה מפורשת, או נכוון את ההגירה מראש עם cudaMemPrefetchAsync. את שיקולי הביצועים האלה נעמיק בפרק על ניהול הזיכרון; כאן העיקר הוא לזהות את nvidia-uvm.ko כרכיב דרייבר נבדל עם קובץ device משלו.

ה-JIT של הדרייבר - PTX to SASS at load

לדרייבר יש תפקיד נוסף שקל לשכוח שהוא שם: הוא גם compiler. בשיעורי פרק 4 ראינו ש-fat binary יכול לשאת PTX (שפת ביניים וירטואלית) ו/או SASS (אסמבלי מקורי לכל גרסת SM). מה קורה כשהיישום מנסה להריץ kernel שעבורו יש בקובץ רק PTX, ואין SASS מתאים ל-GPU הפיזי? כאן נכנס ה-JIT (Just-In-Time) של הדרייבר.

בזמן טעינת המודול - cuModuleLoad בשפת ה-Driver API, או אוטומטית בהפעלת ה-kernel הראשונה תחת ה-Runtime API - libcuda.so לוקחת את ה-PTX המוטמע, קוראת ל-ptxas הפנימי שלה, ומהדרת אותו ל-SASS של ה-GPU הפיזי הנוכחי. זו בדיוק הסיבה שהטמעת PTX נותנת תאימות קדימה: קובץ שהודר עם PTX של compute_90 יוכל לרוץ בעתיד על Blackwell גם בלי compilation מחדש, כי הדרייבר יַהְדֵּר את ה-PTX ל-SASS של הכרטיס החדש בזמן הריצה.

שימו לב לאבחנה החשובה, שכבר פגשנו בשיעור 0.4: ה-JIT הזה הוא של ה-hardware driver, לא של nvcc. כש-nvcc רץ, מזמן סיים; ה-JIT קורה בזמן ריצה, בתוך הדרייבר, על המכונה של המשתמש.

מכיוון שcompilation ב-JIT עולה זמן בהפעלה, הדרייבר שומר את התוצאה ב-cache על הדיסק. ברירת המחדל היא התיקייה ~/.nv/ComputeCache, ואפשר לשלוט בה במשתני סביבה:

export CUDA_CACHE_DISABLE=1        # disables the cache entirely
export CUDA_CACHE_PATH=/tmp/mycache  # alternate location
export CUDA_FORCE_PTX_JIT=1        # forces JIT from PTX even if matching SASS exists

אפשר לראות את זה בפעולה. נהדר תוכנית עם PTX בלבד (בלי SASS), ואז נריץ אותה ונצפה ב-cache מתמלא:

nvcc -gencode arch=compute_90,code=compute_90 vecadd.cu -o vecadd_ptx  # PTX only
rm -rf ~/.nv/ComputeCache
./vecadd_ptx                                                            # triggers JIT
du -sh ~/.nv/ComputeCache

128K    /home/user/.nv/ComputeCache

ה-cache התמלא כי הדרייבר הידר את ה-PTX ל-SASS של sm_90 ושמר את התוצאה. בהרצה הבאה של אותו בינארי, הדרייבר ימצא את ה-SASS מוכן ב-cache ולא יהדר שוב - אלא אם החלפתם GPU או שדרגתם דרייבר, ואז ה-cache יתבטל ויהודר מחדש.

דרייבר מול Toolkit - driver vs toolkit

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

	הדרייבר - driver	ערכת הכלים - CUDA Toolkit
מה נותן	הרצת GPU בזמן ריצה	פיתוח וcompilation
רכיבי מפתח	nvidia.ko + libcuda.so	nvcc + libcudart.so
רמת ה-API	Driver API (cu*)	Runtime API (cuda*)
התקנה	חבילת דרייבר של NVIDIA	חבילת CUDA Toolkit
נבדק עם	nvidia-smi	nvcc --version

הדרייבר מספק את החצי שמריץ בפועל - מודול הליבה שנוגע בחומרה, וספריית ה-Driver API. ה-Toolkit מספק את החצי שבונה - הcompiler nvcc וספריית ה-Runtime API שרוב הקוד קורא לה. אפשר להתקין דרייבר בלי Toolkit (מכונת הרצה גרידא) ואפשר, פחות שכיח, ההפך.

לכל אחד יש פקודת גרסה משלו. הדרייבר:

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 |
|   0  NVIDIA H100 80GB HBM3  On  | 00000000:17:00.0 Off |                    0 |
+-----------------------------------------------------------------------------+

וה-Toolkit:

nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver
Cuda compilation tools, release 12.4, V12.4.131

המלכודת שכולם נופלים בה. בפינה הימנית-עליונה של פלט nvidia-smi כתוב CUDA Version: 12.4. זה אינו אומר שמותקן Toolkit של 12.4, ואף לא שיש בכלל nvcc על המכונה. המספר הזה הוא גרסת ה-CUDA המקסימלית שהדרייבר הזה יודע לתמוך בה. הוא נגזר מגרסת הדרייבר, לא מהcompiler. בהחלט ייתכן ש-nvidia-smi מראה 12.4 בזמן ש-nvcc --version מראה 12.2 - זה תקין לגמרי ואומר שה-Toolkit ישן מהמקסימום שהדרייבר תומך בו.

אפשר לשאול את שתי הגרסאות גם מתוך התוכנית, ולראות שהן שתי מספרות שונות:

#include <cstdio>
#include <cuda_runtime.h>

#define CUDA_CHECK(call)                                                    \
    do {                                                                    \
        cudaError_t err_ = (call);                                          \
        if (err_ != cudaSuccess) {                                          \
            fprintf(stderr, "CUDA error %s at %s:%d\n",                     \
                    cudaGetErrorString(err_), __FILE__, __LINE__);          \
            exit(EXIT_FAILURE);                                             \
        }                                                                   \
    } while (0)

int main() {
    int driverVer = 0, runtimeVer = 0;
    CUDA_CHECK(cudaDriverGetVersion(&driverVer));    // the maximum the driver supports
    CUDA_CHECK(cudaRuntimeGetVersion(&runtimeVer));  // the Toolkit version we compiled against
    printf("Driver  supports CUDA up to: %d.%d\n", driverVer / 1000, (driverVer % 1000) / 10);
    printf("Runtime (toolkit) version:   %d.%d\n", runtimeVer / 1000, (runtimeVer % 1000) / 10);
    return 0;
}

Driver  supports CUDA up to: 12.4
Runtime (toolkit) version:   12.2

שתי הפונקציות מחזירות מספר שלם בפורמט major*1000 + minor*10 (כך ש-12040 הוא 12.4). cudaDriverGetVersion מדווחת על יכולת הדרייבר, ו-cudaRuntimeGetVersion על ה-Toolkit שקומפלנו מולו. מאקרו ה-CUDA_CHECK שכאן הוא בדיוק זה שנשתמש בו לאורך הקורס: הוא עוטף כל קריאת Runtime API, בודק את ה-cudaError_t המוחזר, ומדפיס קובץ, שורה והודעה קריאה אם משהו נכשל - בלעדיו שגיאות נבלעות בשקט.

תאימות קדימה, ומדוע היא חשובה. הכלל: יישום שהודר מול Driver API ישן ירוץ על מערכת עם libcuda.so חדש יותר. משמעותית לפריסה: אתם בונים בינארי פעם אחת, והוא ממשיך לעבוד כשמנהל המערכת משדרג את הדרייבר. מאז CUDA 11 קיימת גם תאימות minor רחבה יותר, שמאפשרת לבינארי שקומפל עם Toolkit של 12.x לרוץ מעל דרייבר שתומך ב-12.0 ומעלה בלי compilation מחדש. את הפרטים המדויקים של מטריצת התאימות נראה כשנדבר על פריסה; לעת עתה זכרו את הכיוון: אפליקציה ישנה על דרייבר חדש - כן; אפליקציה חדשה שדורשת יכולות שהדרייבר הישן לא מכיר - לא בהכרח.

סיכום

• הדרייבר של NVIDIA הוא המתווך היחיד אל ה-GPU: כל אינטראקציה עם החומרה נארזת כבקשה, חוצה אל הליבה, ומבוצעת בתוך nvidia.ko - הרכיב היחיד המורשה לגעת בסיליקון.
• הדרייבר נחתך סביב גבול ה-user/kernel: libcuda.so במרחב המשתמש (לא מיוחסת) ו-nvidia.ko במרחב הליבה (מיוחסת, ring 0) - שני קבצים נפרדים ששניהם יחד הם "הדרייבר".
• מודול הליבה nvidia.ko הוא kernel object בפורמט ELF, נטען עם modprobe, נבדק עם lsmod ו-modinfo, וכותב הודעות NVRM ל-dmesg; לצדו נטענים nvidia-uvm.ko, nvidia-modeset.ko ו-nvidia-drm.ko.
• חצי הליבה שוחרר כקוד פתוח (NVIDIA/open-gpu-kernel-modules, רישיון Dual MIT/GPL) והוא ברירת המחדל ל-H100; אך libcuda.so במרחב המשתמש נשארה סגורה (עם חלופות כמו LibreCuda ו-tinygrad).
• הספרייה libcuda.so מממשת את ה-Driver API (cu*), מקושרת תמיד דינמית, ומגיעה עם הדרייבר ולא עם ה-Toolkit; ה-soname libcuda.so.1 קבוע על פני גרסאות ומאפשר תאימות קדימה, והיעדרה מסמן דרייבר שלא הותקן.
• נתיב הבקשה הוא קוד משתמש -> libcudart.so -> libcuda.so -> ioctl על קובצי /dev/nvidia* -> nvidia.ko -> חומרה; ה-ioctl (magic 0x46 = 'F') הוא הpipe הצר של כל האינטראקציה, ואפשר לראותו עם strace.
• רכיב ה-UVM (nvidia-uvm.ko, קובץ /dev/nvidia-uvm) מממש מרחב כתובות מאוחד ל-host ול-device מאחורי cudaMallocManaged, ומהגר דפים על page fault - נוחות שעולה ב-latency לעומת הקצאה והעתקה מפורשות.
• הדרייבר גם compiler: ה-JIT שלו מתרגם PTX ל-SASS בזמן טעינה (זה ה-hardware driver, לא nvcc), מה שנותן תאימות קדימה, ושומר את התוצאה ב-~/.nv/ComputeCache (נשלט ב-CUDA_CACHE_DISABLE, CUDA_FORCE_PTX_JIT).
• הדרייבר (nvidia.ko + libcuda.so, נבדק ב-nvidia-smi) וה-Toolkit (nvcc + libcudart.so, נבדק ב-nvcc --version) הן שתי חבילות נפרדות; ה-CUDA Version שב-nvidia-smi הוא המקסימום שהדרייבר תומך בו, לא גרסת ה-Toolkit המותקנת.