0.4 מפת שרשרת התוכנה הרצאה

בשיעורים הקודמים של פרק ההקדמה הכרנו את ה-GPU (מעבד גרפי / כרטיס מסך) כמאיץ מקבילי, ראינו למה הוא שונה במהותו מ-CPU (מעבד), והתוודענו לשבב שילווה אותנו לאורך כל הקורס - ה-NVIDIA H100. אבל עד עכשיו דיברנו בעיקר על החומרה. השיעור הזה סוגר את הפרק בעזרת מפה: המסע המלא מהשורה שאתם כותבים ב-CUDA C++ ועד להוראה הבודדת שרץ עליה טרנזיסטור בתוך ה-SM (מעבד רב-זרמי - Streaming Multiprocessor). זו הרצאה ניווטית. לא נלמד בה לכתוב kernel (קרנל - פונקציית ה-GPU) מהיר, אלא נבנה את שלד המושגים שכל שאר הקורס יתלה עליו את הבשר. בסוף השיעור תדעו להצביע על כל רכיב בשרשרת - CUDA C++, nvcc, PTX, SASS, libcudart.so, libcuda.so, nvidia.ko - ולהגיד בדיוק מה תפקידו, מי מעליו ומי מתחתיו. את הפרטים המלאים על הcompiler ושתי שפות הביניים נפרוש בפרק 4, ואת מחסנית ה-host במלואה בפרק 5. כאן אנחנו מציירים את התמונה הגדולה.

שלוש המשמעויות של המילה CUDA - CUDA¶

לפני שנצייר מפה, צריך לפרק מוקש מושגי אחד ולתמיד. המילה CUDA (ראשי תיבות של Compute Unified Device Architecture) עמוסה מדי, והיא מציינת שלושה דברים שונים לגמרי. חוסר הבחנה ביניהם הוא מקור בלבול קבוע, אז נעצור עליהם כאן:

ארכיטקטורת device - device architecture. במובן הזה CUDA היא תכן החומרה של ה-GPU עצמו - האופן שבו בנויים ה-SM, ה-CUDA Core (ליבת CUDA), ה-Tensor Core (ליבת טנזור) וההיררכיה שסביבם. זו המשמעות המקורית של שלוש האותיות האחרונות - Device Architecture.
מודל תכנות מקבילי - programming model. במובן הזה CUDA היא הפרדיגמה: היררכיית ה-threads (threads), היררכיית הזיכרון, ה-kernel ש"נקרא פעם אחת ורץ פעמים רבות". זו הפשטה מופשטת שאינה תלויה בשפה מסוימת.
פלטפורמת תוכנה - software platform. במובן הזה CUDA היא אוסף הכלים הקונקרטי שמרחיב שפות כמו C++ עם מודל התכנות הזה - הcompiler nvcc, ה-Runtime API, ה-Driver API, וספריות כמו cuBLAS ו-cuDNN שיושבות מעליהם.

השיעור הזה עוסק כמעט כולו במשמעות השלישית - הפלטפורמה, כלומר השרשרת שהופכת קוד לחומרה שרצה. שימו לב שהמודל הרעיוני (משמעות 2) הוא מימוש (implementation) מעל הארכיטקטורה (משמעות 1), ופלטפורמת התוכנה (משמעות 3) היא מימוש קונקרטי של מודל התכנות. כשמישהו אומר לכם "זה רץ ב-CUDA", שאלו את עצמכם: הוא מדבר על החומרה, על הפרדיגמה, או על הכלים.

הערה היסטורית ששווה להכיר: כמעט כל הניסוחים, הדיאגרמות ואפילו כמה מהמשפטים בתיעוד הרשמי של NVIDIA נשענים על מאמר היסוד Lindholm et al. משנת 2008. הוא מקור ראשוני טוב לצטט ממנו לאורך הקורס.

נתיב הבנייה - build path¶

נתיב הבנייה הוא מה שקורה לפני ההרצה, על ה-host, בזמן compilation (compile time). המקור שאתם כותבים הוא CUDA C++ - הרחבה של שפת C++ הרגילה. ההרחבה מוסיפה בדיוק את מילות המפתח שהופכות C++ רגילה לתכנות GPU:

#include <cstdio>
#include <cstdlib>

// error-checking macro - we'll explain it shortly
#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)

// kernel: a function marked __global__, returns void, receives pointers
__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // unique global index for each thread
    if (i < n)                                        // bounds guard against over-launch
        c[i] = a[i] + b[i];                           // write the result through a pointer
}

int main() {
    const int n = 1 << 20;                            // a million elements
    size_t bytes = n * sizeof(float);
    float *da, *db, *dc;
    CUDA_CHECK(cudaMalloc(&da, bytes));
    CUDA_CHECK(cudaMalloc(&db, bytes));
    CUDA_CHECK(cudaMalloc(&dc, bytes));

    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;  // round up
    vecAdd<<<numBlocks, blockSize>>>(da, db, dc, n);   // launch from the host

    CUDA_CHECK(cudaGetLastError());                    // catches launch errors
    CUDA_CHECK(cudaDeviceSynchronize());               // waits for the kernel to finish
    CUDA_CHECK(cudaFree(da));
    CUDA_CHECK(cudaFree(db));
    CUDA_CHECK(cudaFree(dc));
    return 0;
}

בקוד הזה מרוכזות כל התוספות המרכזיות ש-CUDA C++ מוסיפה ל-C++ רגילה:

מילת המפתח __global__ מגדירה kernel. שימו לב לשני עקרונות: ה-kernel תמיד מחזיר void, והוא מתקשר עם העולם דרך מצביעים (pointers) - הוא כותב תוצאות לזיכרון, לא "מחזיר" ערך. טעות מתחילים נפוצה היא לצפות ש-kernel יחזיר מספר. מלכוד שמות חשוב: __global__ אינו קשור ל-global memory (זיכרון גלובלי); הוא מציין פונקציה שמופעלת מה-host ורצה על ה-device. אחד ממעצבי CUDA המוקדמים אמר שהשם הזה נבחר "למקסימום בלבול של מפתחים", אז שימו לב.
תחביר ההפעלה <<<>>> בין שם הפונקציה לארגומנטים קובע את מימדי ה-grid (רשת) של ה-blocks (בלוקים). זהו מבנה של קוד ה-host - הוא מופיע בקוד ה-CPU כדי להפעיל קוד device.
מילת המפתח __shared__, האינטרינזיק __syncthreads(), והמשתנים המובנים blockDim, threadIdx, blockIdx זמינים בתוך קוד ה-device. הם מאפשרים לכל thread לחשב על איזה נתח נתונים הוא אחראי, ולשתף פעולה בתוך ה-block. בקוד למעלה השתמשנו רק ב-blockIdx, blockDim ו-threadIdx לחישוב האינדקס; את __shared__ ואת ה-barrier נראה כשנכתוב kernel אמיתי מהיר.

על מאקרו בדיקת השגיאות CUDA_CHECK. כל קריאת Runtime API מחזירה cudaError_t. אם לא בודקים אותה, שגיאות נבלעות בשקט וה-kernel פשוט לא רץ בלי להתלונן. המאקרו עוטף כל קריאה, ומדפיס קובץ, שורה והודעה קריאה אם משהו נכשל. שימו לב ל-cudaGetLastError() מיד אחרי ההפעלה: הפעלת kernel בעצמה אינה מחזירה קוד שגיאה דרך הערך המוחזר, אז זו הדרך לתפוס שגיאות הפעלה (למשל grid לא חוקי). נשתמש במאקרו הזה לאורך כל הקורס.

הcompilation עצמו הוא מפוצל. הפקודה שבונה את הקובץ היא:

nvcc -arch=sm_90 vecadd.cu -o vecadd

מנהל-הcompilation nvcc (compiler-מנהל של CUDA) אינו compiler מונוליטי אחד אלא toolchain - אורקסטרטור שמפעיל כמה כלים מתחתיו. הוא מפריד את קוד ה-device מקוד ה-host: את קוד ה-host הוא מוסר לcompiler המערכת (למשל gcc), ואת קוד ה-device הוא מהדר בעצמו בהמשך ל-PTX ו/או ל-SASS. התוצאה היא קובץ הרצה יחיד.

התוצר: קובץ שמן - fat binary. קובץ ההרצה תואם ל-host ABI (Application Binary Interface - interface בינארי יישומי), כלומר הוא קובץ הרצה רגיל של ה-CPU - אבל הוא נושא בתוכו את קוד ה-GPU מוטמע. הקוד המוטמע הזה יכול להיות PTX, SASS, או שניהם. מכיוון שזה קובץ host רגיל, אפשר להציץ בו עם כלים סטנדרטיים:

readelf -S vecadd | grep nv_fatbin

  [23] .nv_fatbin        PROGBITS         0000000000019a80  00019a80

הפלט מראה מקטע (section) בשם .nv_fatbin בתוך קובץ ELF רגיל לחלוטין - שם יושב ה-fat binary. לניתוח מודע-CUDA נשתמש ב-CUDA Binary Utilities (כלי-העזר לקבצים בינאריים של CUDA), ובראשם cuobjdump:

cuobjdump --list-elf vecadd     # which architectures are embedded
cuobjdump -ptx vecadd           # extracts the PTX
cuobjdump -sass vecadd          # extracts the SASS

בברירת המחדל, -arch=sm_90 הוא קיצור ל--gencode arch=compute_90,code=sm_90, שמטמיע SASS עבור sm_90 בלבד - בלי PTX. כדי להטמיע גם PTX (לתאימות עתידית) וגם SASS, מבקשים במפורש את שניהם:

nvcc -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     vecadd.cu -o vecadd

השורה הראשונה מטמיעה SASS ל-Hopper; השנייה מטמיעה PTX של compute_90. מיד נבין מה ההבדל בין השניים ולמה חשוב לשלוט בו.

שתי שפות הביניים - PTX מול SASS¶

קוד ה-device עובר דרך שתי ייצוגים שקל מאוד לבלבל ביניהם. שניהם קוד GPU, אבל הם ברמות שונות לגמרי, וההבחנה ביניהם היא אחד הרעיונות המרכזיים בכל שרשרת הכלים.

PTX - הרצת שרשורים מקבילית - Parallel Thread eXecution. זו שפת ביניים וירטואלית (virtual ISA / IR - Intermediate Representation), בדיוק במובן ש-LLVM-IR הוא ייצוג ביניים. NVIDIA מתארת אותה גם כ"מכונה וירטואלית" וגם כ"ארכיטקטורת סט הוראות". הרעיון: מי שמייצר PTX יכול לסמוך שהתוכנית תרוץ באותה סמנטיקה על הרבה מכונות פיזיות שונות - כולל מכונות שעוד לא קיימות. שלא כמו x86_64 או aarch64, PTX אינו קוד מכונה אמיתי אלא IR שעובר compilation בזמן ריצה (JIT - Just-In-Time) על ידי דרייבר ה-GPU אל ה-SASS של המכשיר הספציפי. הנה שלוש שורות PTX אמיתיות:

.reg .f32 %f<7>;
fma.rn.f32 %f5, %f4, %f3, 0f3FC00000;
mov.u32 %r1, %ctaid.x;

השורה הראשונה היא הנחיה לcompiler: ה-kernel צורך שבעה registers (רגיסטרים / אוגרים) מסוג 32-bit floating point. הרגיסטרים מוקצים דינמית לקבוצות של threads (warps - קבוצות של 32 threads) מתוך קובץ האוגרים של ה-SM. השורה השנייה היא fused multiply-add: הכפל את f3 ב-f4, הוסף את הקבוע (הערך 0f3FC00000 הוא 1.5 בייצוג IEEE 754), ואחסן ב-f5. הסיומת rn קובעת עיגול "round even" של IEEE 754. השורה השלישית מעתיקה את ערך ה-x של אינדקס ה-CTA (Cooperative Thread Array - המקבילה ברמת ה-ISA ל-block) לרגיסטר.

SASS - אסמבלר זורם - Streaming ASSembler. זו שפת האסמבלי המקורית של ה-GPU, הרמה הנמוכה ביותר שבה עדיין אפשר לכתוב קוד קריא לאדם. ה"Streaming" בשם מתייחס ל-SM שאת ההוראות שלהם השפה מתארת. שלא כמו PTX, SASS מקושר לגרסת SM פיזית מסוימת - הוא אינו נייד בין דורות. הנה SASS אמיתי מ-SM90a (הארכיטקטורה של Hopper, כלומר H100):

FFMA R0, R7, R0, 1.5 ;
S2UR UR4, SR_CTAID.X ;

השורה הראשונה היא Fused Floating-point Multiply-Add על רגיסטרים כלליים (R). השנייה מעתיקה את ערך ה-x של אינדקס ה-CTA מ-Special Register (רגיסטר מיוחד, SR) אל Uniform Register 4 (רגיסטר אחיד - UR, ערך זהה לכל ה-warp). שימו לב שאותה פעולת fma מתוך ה-PTX הפכה כאן ל-FFMA קונקרטית - זה בדיוק מה שה-JIT של הדרייבר עושה.

ההשוואה בטבלה:

מאפיין	PTX	SASS
רמה	שפת ביניים וירטואלית (IR)	אסמבלי מקורי של החומרה
אנלוגיה	LLVM-IR	קוד מכונה של x86
מקושר לגרסת	compute capability (compute_XX)	גרסת SM (sm_XX)
ניידות	קדימה-תואם על פני דורות	לא נייד, לא יציב בין דורות
מי מייצר	nvcc (וגם NVRTC)	ptxas, או ה-JIT של הדרייבר מ-PTX
מתי הופך ל-SASS	בזמן ריצה (JIT) אם מוטמע כ-PTX	כבר הוא SASS
תיעוד הקידוד	מתועד רשמית	הסמנטיקה לא מוגדרת, הקידוד הבינארי לא מתועד

מסקנה מעשית: PTX הוא ה"מותן הצר" (narrow waist) שמפריד בין עולם התוכנה לעולם החומרה - interface יציב שמעליו compilers ומתחתיו מכשירים. מטרגטים אליו לצורך עמידות. יורדים ל-inline PTX בכתב יד רק עבור תכונות שהcompiler עדיין לא יודע לפלוט - נכון ל-2025, זו הדרך היחידה לנצל תכונות ספציפיות ל-Hopper כמו ההוראות wgmma (warpgroup matrix-multiply-accumulate) ו-tma (Tensor Memory Accelerator), כמו ב-Flash Attention 3. את SASS כמעט אף אחד לא כותב ביד; מסתכלים עליו בזמן profiling ואז עורכים את ה-CUDA C++ או את ה-PTX. את כל זה נעמיק בפרק 4.

יכולת החישוב שמחברת קוד לחומרה - compute capability¶

מה מחליט אילו GPU פיזיים יכולים להריץ PTX או SASS מסוימים? מספר גרסה בשם compute capability. זו מערכת הגרסאות שמפשיטה את פרטי ה-GPU הפיזי מסט ההוראות ומהcompiler - שכבת ההפרדה בין הקוד לסיליקון.

הפורמט הוא major.minor, למשל 7.0 או 9.0. הנה המיפוי לדורות שמלווים אותנו בקורס:

compute capability	דור - generation	דוגמה
7.0	Volta	V100
7.5	Turing	T4
8.0	Ampere	A100
8.9	Ada Lovelace	L4 / RTX 4090
9.0	Hopper	H100
10.0	Blackwell	B200

ה-H100 שלנו הוא compute capability 9.0. אפשר לשאול את הכרטיס ישירות:

nvidia-smi --query-gpu=name,compute_cap --format=csv

name, compute_cap
NVIDIA H100 80GB HBM3, 9.0

וירטואלי מול פיזי - זו הבחנה אמיתית. תיעוד nvcc מכנה את ה-compute capability "ארכיטקטורת GPU וירטואלית" (compute_XX), בניגוד ל"ארכיטקטורת GPU פיזית" (sm_XX). לכן compute_90 הוא הטרגט הווירטואלי (ומייצר PTX) ו-sm_90 הוא הטרגט הפיזי (ומייצר SASS). בברירת מחדל, כשמטרגטים compute capability, הcompiler מייצר גם PTX לארכיטקטורה הווירטואלית וגם SASS מותאם ל-SM הפיזי המתאים.

תאימות קדימה - forward compatibility. גרסאות סטנדרטיות מבטיחות ש-PTX ישן ירוץ על GPU חדש יותר, לפי "מודל שכבות הבצל" (onion model). המנגנון הוא ה-JIT: ה-PTX המוטמע בקובץ עובר compilation לזמן ריצה ל-SASS של המכשיר החדש. זו הסיבה שכדאי להטמיע PTX ולא רק SASS - קובץ שהודר עם PTX של compute_90 יוכל לרוץ בעתיד על Blackwell גם בלי compilation מחדש, בעוד SASS של sm_90 מקובע ל-Hopper.

זהירות: ההבטחה כבר לא אחידה. Hopper הכניסה סיומת a (למשל 9.0a) ו-Blackwell סיומת f (למשל 10.0f) לתכונות שסוטות ממודל הבצל. עבור a אין הבטחת תאימות עתידית אפילו בתוך אותה major; עבור f יש תאימות על פני minor אך לא על פני major. ההוראות wgmma ו-tma שהזכרנו הן בדיוק תכונות sm_90a כאלה. את הכללים המדויקים לכל גרסה נראה בפרק 4; לעת עתה, זכרו את משוואת היסוד: compute capability = וירטואלי (compute_XX, גוזר PTX), SM version = פיזי (sm_XX, גוזר SASS).

נתיב ההרצה - run path¶

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

שכבה 1 - CUDA Runtime API. זהו הinterface הגבוה יותר, והנפוץ ביותר. הפונקציות שלו נושאות את הקידומת cuda (למשל cudaMalloc, cudaMemcpy, cudaDeviceSynchronize), והתחביר <<<>>> שראינו מקומפל לקריאות שלו. הוא ממומש בספרייה libcudart.so, שמגיעה עם ה-CUDA Toolkit. את ה-vecAdd שלנו כתבנו כולו מול השכבה הזו.

שכבה 2 - CUDA Driver API. זהו הinterface ברמה נמוכה יותר, שנותן שליטה עדינה יותר (ניהול context מפורש, טעינת modules, טעינת PTX ב-JIT ידני). הפונקציות שלו נושאות את הקידומת cu (למשל cuMemAlloc, cuInit, cuLaunchKernel), והוא ממומש בספרייה libcuda.so. מעט מאוד תוכניות נכתבות ישירות מולו - רוב העולם משתמש ב-Runtime API, וה-Runtime API עצמו יושב מעל ה-Driver API.

שכבה 3 - nvidia.ko. זהו מודול הליבה (kernel module) של הדרייבר בלינוקס. הוא רץ במצב מיוחס (privileged mode / ring 0) ומדבר ישירות עם חומרת ה-GPU. הוא החלק היחיד בכל השרשרת שמותר לו לגעת בסיליקון הפיזי; קוד במרחב המשתמש חייב לחצות את גבול ה-syscall/ioctl אל תוך המודול הזה כדי להגיע לחומרה. חלקו שוחרר כקוד פתוח תחת השם Linux Open GPU Kernel Module.

שכבה 4 - חומרת ה-GPU (device). כאן, סוף סוף, רץ ה-SASS על ה-SM.

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

מאפיין	CUDA Runtime API	CUDA Driver API
קידומת פונקציות	`cuda*` (למשל `cudaMalloc`)	`cu*` (למשל `cuMemAlloc`)
רמה	גבוהה, נוחה, נפוצה	נמוכה, שליטה עדינה
ספרייה ממַמֶשת	`libcudart.so`	`libcuda.so`
מגיעה עם	CUDA Toolkit	הדרייבר של NVIDIA
קוד פתוח	חלקית	סגור (יש ניסיונות חלופיים: LibreCuda, tinygrad)

מלכוד שמות שחייבים לזכור: libcuda.so (ה-Driver API, חלק מהדרייבר) הוא לא libcudart.so (ה-Runtime API, חלק ה-toolkit - שימו לב ל-rt). שני קבצים שונים בשתי שכבות שונות. אבחון מעשי: אם libcuda.so חסר מהמערכת, המשמעות בדרך כלל היא שהדרייבר של ה-GPU לא הותקן כראוי. עוד עובדה תפעולית: ה-Driver API קדימה-תואם - יישום שהודר מול גרסה ישנה ירוץ על מערכת עם libcuda.so חדש יותר. את כל המחסנית הזו נפרק לרכיביה בפרק 5.

הcompiler-מנהל אינו מנהל-device - compiler driver vs hardware driver¶

לפני שנרכיב את המפה השלמה, נסלק עוד מוקש - הפעם מוקש של מילים. שתי מילים אנגליות שונות מתורגמות לעברית ל"מנהל", והמילה האנגלית "driver" מופיעה בשני מקומות שאין ביניהם שום קשר:

הcompiler-מנהל - compiler driver. זהו nvcc, וזה גם gcc. תפקידו לארכֵּן (orchestrate) את תהליך הcompilation - להפעיל את הכלים שמתחתיו בסדר הנכון. הוא כלי של זמן compilation, על ה-host, ואין לו שום קשר לחומרה.
מנהל-device - hardware driver. זהו הדרייבר של ה-GPU של NVIDIA - שכבות libcuda.so ו-nvidia.ko שראינו בנתיב ההרצה. תפקידו לתווך בין ה-host לבין ה-GPU בזמן ריצה.

אותה מילה, "driver", שני דברים שונים לחלוטין. התיעוד הרשמי של NVIDIA מזהיר על כך במפורש: אל תבלבלו compiler drivers עם hardware drivers. כשאתם קוראים "the driver compiles the PTX", עצרו ושאלו איזה driver: את PTX מהדר ה-hardware driver (nvidia.ko דרך ה-JIT) - לא nvcc. וכשאתם קוראים "nvcc drives the compilation", זה ה-compiler driver, שאינו נוגע בכרטיס. הבחנה זו תמנע מכם חצי מבלבולי הקורס.

מפת השכבות המלאה - the full stack¶

עכשיו אפשר להרכיב הכל למפה אחת שנחזור אליה שוב ושוב. מלמעלה (הקוד שאתם כותבים) עד למטה (החומרה שרצה):

                 Build path                        Run path
              (compile time, host)            (runtime, host -> device)

  +-----------------------------------------------------------------+
  |  CUDA C++  (__global__, <<<>>>, __shared__, __syncthreads,      |   <- the code you write
  |            blockDim / threadIdx / blockIdx)                      |
  +-----------------------------------------------------------------+
           |  nvcc  (compiler driver)                                  host code -> gcc
           |  splits: host -> gcc  |  device -> PTX / SASS
           v
  +-----------------------------------------------------------------+
  |  fat binary  (ELF file matching host ABI, .nv_fatbin section)   |   readelf / cuobjdump
  |    +---------------------+   +-------------------------------+   |
  |    |  PTX  (compute_90)  |   |  SASS  (sm_90)                |   |
  |    |  virtual ISA / IR   |   |  native, per-SM assembly      |   |
  |    +---------------------+   +-------------------------------+   |
  +-----------------------------------------------------------------+
                                        |
   ======================= boundary between build and run ===========
                                        |
                                        v
  +-----------------------------------------------------------------+
  |  the application                                                 |   the running host code
  +-----------------------------------------------------------------+
           |  calls the API, "in order":
           v
  +-----------------------------------------------------------------+
  |  CUDA Runtime API   (cuda*,  libcudart.so)   <- high-level, common |   user mode
  +-----------------------------------------------------------------+
           |
           v
  +-----------------------------------------------------------------+
  |  CUDA Driver API    (cu*,    libcuda.so)     <- low-level         |   user mode
  +-----------------------------------------------------------------+
   ---------------------- user / kernel boundary ---------------------
           |  (syscall / ioctl)
           v
  +-----------------------------------------------------------------+
  |  nvidia.ko   (kernel module, privileged mode)                   |   kernel mode
  |  this is also where the JIT happens: PTX -> SASS if only PTX    |
  |  was embedded                                                    |
  +-----------------------------------------------------------------+
           |  touches the hardware directly
           v
  +-----------------------------------------------------------------+
  |  GPU hardware (device):  132 SMs of the H100, CUDA Cores,       |   silicon
  |  Tensor Cores, register file, shared memory, HBM3 80 GiB        |
  +-----------------------------------------------------------------+

בואו נלך על מסלול קונקרטי אחד מקצה לקצה, עם ה-vecAdd שלנו על H100:

אתם כותבים vecadd.cu ב-CUDA C++.
הפקודה nvcc -arch=sm_90 vecadd.cu -o vecadd מפעילה את nvcc (compiler driver): קוד ה-host הולך ל-gcc, קוד ה-kernel הופך ל-SASS של sm_90 (ואם ביקשתם, גם ל-PTX של compute_90).
התוצר vecadd הוא fat binary - קובץ ELF רגיל עם מקטע .nv_fatbin. readelf יראה לכם את המקטע, cuobjdump -sass יראה את ה-FFMA.
אתם מריצים ./vecadd. הקריאה cudaMalloc נכנסת ל-libcudart.so (Runtime API), שקוראת בתורה ל-libcuda.so (Driver API).
libcuda.so חוצה אל nvidia.ko במצב מיוחס. אם הקובץ הכיל רק PTX, כאן מתבצע ה-JIT ל-SASS של sm_90.
nvidia.ko משגר את ה-SASS ל-GPU. ה-grid של 4096 blocks בגודל 256 (סך 1,048,576 threads) מתפרש על 132 ה-SM של H100, וכל thread מבצע חיבור אחד.

זו המפה. כל שיעור בהמשך הקורס ממקד זכוכית מגדלת על תא אחד בטבלה הזו - פרק 4 על PTX, SASS ו-nvcc, פרק 5 על ה-Runtime, ה-Driver ו-nvidia.ko, והפרקים על החומרה על השורה התחתונה של הדיאגרמה.

סיכום¶

המילה CUDA עמוסה בשלוש משמעויות שיש להפריד: ארכיטקטורת device (חומרה), מודל תכנות (פרדיגמה), ופלטפורמת תוכנה (הכלים) - השיעור עסק בעיקר בשלישית.
נתיב הבנייה מתחיל ב-CUDA C++ עם __global__, <<<>>>, __shared__, __syncthreads, blockDim ו-threadIdx, ועובר דרך nvcc שמפצל את הקוד: host ל-gcc, device ל-PTX ו/או SASS.
התוצר הוא fat binary - קובץ ELF תואם host ABI עם מקטע .nv_fatbin, שאפשר לבדוק עם readelf (גנרי) או cuobjdump (מודע-CUDA).
קרנל תמיד מחזיר void ומתקשר דרך מצביעים; __global__ אינו קשור ל-global memory אלא מסמן פונקציה שמופעלת מה-host ורצה על ה-device.
PTX הוא שפת ביניים וירטואלית (כמו LLVM-IR), מגורסֶת לפי compute capability, קדימה-תואמת, ועוברת JIT על ידי הדרייבר; SASS הוא אסמבלי מקורי לכל גרסת SM, לא יציב וקידודו לא מתועד.
compute capability היא מערכת הגרסאות שמחברת PTX/SASS ל-GPU פיזיים: compute_XX וירטואלי (גוזר PTX) מול sm_XX פיזי (גוזר SASS); ה-H100 הוא 9.0.
נתיב ההרצה יורד מהיישום ל-Runtime API (cuda*, libcudart.so, גבוה ונפוץ), ל-Driver API (cu*, libcuda.so, נמוך), ל-nvidia.ko (מודול ליבה מיוחס), ולבסוף לחומרה.
אל תבלבלו בין compiler driver (nvcc, gcc - מארכן compilation) לבין hardware driver (הדרייבר של NVIDIA - libcuda.so ו-nvidia.ko): אותה מילה "driver", שני דברים שונים לגמרי.
אבחון מהיר: libcuda.so חסר משמעו שהדרייבר לא הותקן; libcuda.so (Driver) הוא לא libcudart.so (Runtime) - שני קבצים בשתי שכבות.
את הפרטים המלאים של הcompiler ושפות הביניים נעמיק בפרק 4, ואת מחסנית ה-host במלואה בפרק 5; המפה בשיעור הזה היא המצפן לשניהם.