4.3 nvcc, nvrtc וההידור הרצאה
בשיעור 4.1 הכרנו את PTX - שפת האסמבלי הווירטואלית והנייחת של NVIDIA - ובשיעור 4.2 ראינו את SASS, האסמבלי האמיתי של החומרה שרץ בפועל על ה-SM, ולמדנו לקרוא אותו. עכשיו נשאלת השאלה המעשית שסוגרת את הפרק: מי מייצר את שני ה-artifacts האלה מתוך קובץ ה-.cu שכתבנו, ואיך הם נארזים לתוך תוכנית אחת שרצה? התשובה היא nvcc - וההפתעה הראשונה, שנפתח בה, היא ש-nvcc בכלל אינו מהדר (compiler) אחד אלא מנהל-הידור (compiler driver), מנצח שמתזמר שרשרת שלמה של כלים קטנים יותר. בשיעור הזה נפרק את השרשרת הזאת חוליה-חוליה: איך nvcc מפצל את הקוד לhost ולdevice, מי מתרגם CUDA C++ ל-PTX ומי מתרגם PTX ל-SASS, מהו ה-"fat binary" שנוצר ואיך מציצים לתוכו, איך שולטים באילו גרסאות קוד נארזות פנימה עם compute_XX ו-sm_XX, איך המאקרו __CUDA_ARCH__ מאפשר קוד שונה לכל דור חומרה, איך מקשרים __device__ בין קבצים נפרדים, ולבסוף - איך nvrtc עושה את כל הcompilation הזה בזמן ריצה, בתוך התהליך שלנו, בלי להריץ nvcc כלל. ה-compute capability שהכרנו בשיעור 2.5 יהיה הציר המרכזי של כל הדיון הזה, כי הוא בדיוק מה שמספרר את הקוד המוטמע.
הcompiler-מנהל nvcc ולא compiler יחיד - the compiler driver¶
הדבר הראשון שחשוב להפנים: nvcc הוא toolchain, כלומר תזמורת של כלים, ולא compiler מונוליטי. כשאנחנו כותבים nvcc mm.cu -o mm, מאחורי הקלעים nvcc מבצע פיצול (splitting) של קובץ המקור לשני זרמים - קוד host (ה-CPU) וקוד device (ה-GPU) - ומעביר כל זרם למסלול הcompilation המתאים לו. את קוד הhost nvcc אינו compiler בעצמו: הוא מוסר אותו לcompiler של מערכת ההפעלה, כלומר ל-gcc (או clang / cl.exe על Windows). את קוד הdevice הוא מעביר דרך שרשרת כלי ה-GPU הפנימיים.
הנה השרשרת המלאה, מהמקור ועד לקובץ ההרצה:
mm.cu (CUDA C++)
│
┌──────┴───────┐ nvcc splits (cudafe++)
▼ ▼
host code device code
│ │
│ ▼
│ cicc CUDA C++ ─► PTX
│ │ (mm.ptx)
│ ▼
│ ptxas PTX ─► SASS (cubin)
│ │ (mm.cubin)
│ ▼
│ fatbinary packs PTX + SASS into fatbin
│ │ (mm.fatbin)
▼ ▼
gcc ◄──── embedded as a byte array in the host code
│
▼
ld / link host ABI executable
│
▼
mm host binary carrying GPU code inside
חמשת הכלים המרכזיים שכדאי לזכור בשמם:
| כלי | תפקיד | קלט -> פלט |
|---|---|---|
cudafe++ |
מפצל את המקור לקוד host ולקוד device | .cu -> קבצי ביניים |
cicc |
הcompiler של קוד הdevice (מבוסס LLVM) | CUDA C++ -> PTX |
ptxas |
ה-assembler האופטימיזציוני של PTX | PTX -> SASS (cubin) |
fatbinary |
אורז את ה-PTX וה-SASS לחבילה אחת | PTX + cubin -> fatbin |
nvlink |
מקשר קוד device בין קבצים (compilation נפרד) | קבצי .o -> device link |
שימו לב לחלוקת העבודה בין שני הכלים החשובים ביותר. cicc הוא זה שהופך את ה-CUDA C++ ל-PTX (זה מה שראינו בשיעור 4.1). ptxas, למרות ששמו רומז ל"assembler" פשוט, הוא בעצם compiler אופטימיזציוני מלא: הוא זה שמבצע הקצאת אוגרים (register allocation), פריסת לולאות, תזמון הוראות ובחירת ההוראות הסופיות של SASS (שיעור 4.2). ה--O3 שאתם מכירים מ-gcc תקף ל-cicc, אבל את מספר האוגרים ואת האופטימיזציות ברמת ה-SASS קובע ptxas - ולכן דגלים כמו -Xptxas -v (הדפסת שימוש באוגרים ו-shared memory) ו---maxrregcount=N הם שער השליטה החשוב ביותר על הביצועים ברמת ה-kernel. הדגל -Xptxas מעביר ארגומנטים ל-ptxas, -Xcicc ל-cicc, ו--Xcompiler לcompiler הhost.
רוצים לראות את השרשרת במו עיניכם? שני דגלים:
nvcc --dryrun mm.cu -o mm # prints all the sub-commands without running them
nvcc --keep mm.cu -o mm # keeps the intermediate files: mm.ptx, mm.cubin, mm.fatbin ...
המלכוד המושגי המרכזי של השיעור (misconception). המילה "driver" מופיעה כאן בשני מובנים שאין ביניהם דבר וחצי דבר. nvcc הוא compiler driver - מנהל-הידור, כלי שמתזמר compilers. ה-NVIDIA GPU Driver (מה שנתקין במערכת, nvidia.ko ו-libcuda, נלמד בפרק 5) הוא hardware driver - תוכנה שמדברת עם החומרה. אותה מילה, שני עולמות. אל תבלבלו: nvcc (וגם gcc) הם compiler drivers; הם רצים בזמן ה-build ולא נוגעים ב-GPU כלל. את התרגום הסופי של PTX ל-SASS בזמן ריצה (JIT) עושה דווקא ה-hardware driver, כפי שנראה בסעיף ה-nvrtc.
הבינארי השמן - the fat binary¶
מה בדיוק nvcc מייצר? קובץ הרצה שמציית ל-host ABI (Application Binary Interface, interface הבינארי של מערכת ההפעלה) - כלומר, מבחינת מערכת ההפעלה זהו קובץ ELF רגיל לחלוטין (על לינוקס), שאפשר להריץ, לקשר ולבדוק עם כל כלי מערכת סטנדרטי. ההבדל היחיד הוא שבתוכו, בסקשן ELF ייעודי, טמון גם קוד ה-GPU. את האריזה הזאת - PTX ו/או SASS ארוזים יחד - מכנים fat binary (fatbin), "בינארי שמן", כי הוא נושא בתוכו קוד לכמה ארכיטקטורות במקביל.
מכיוון שהמעטפת היא ELF תקני, אפשר להסתכל עליה עם הכלים הגנריים של לינוקס. הסקשן שנושא את ה-GPU code נקרא .nv_fatbin:
readelf -S mm | grep -i nv
# .nv_fatbin PROGBITS ... # this is where all the GPU code is hidden
# __nv_relfatbin ...
אבל readelf רואה רק "בלוק בייטים אטום" - הוא לא מבין מה יש בפנים. כדי לפרק את התוכן הפנימי ולראות אילו cubins ואילו גרסאות PTX ארוזות, משתמשים בכלי-העזר לבינאריים של CUDA - CUDA Binary Utilities (הכרנו אותם בשיעור 4.2). הכלי המרכזי הוא cuobjdump, שיודע לפתוח את הבינארי כולו ולהציג את ה-cubins המוטמעים:
cuobjdump -lelf mm # list of the packed cubins (SASS), one per sm_XX
cuobjdump -lptx mm # list of the packed PTX versions, one per compute_XX
cuobjdump -sass mm # disassembly of the SASS itself (what ptxas produced)
cuobjdump -ptx mm # prints the embedded PTX
cuobjdump -res-usage mm # register and shared memory usage per kernel
ה-cuobjdump -lelf הוא בדיוק החלון שמראה לנו את מבנה ה-fatbin: כמה גרסאות SASS יש, לאילו ארכיטקטורות (sm_70, sm_90...), ואילו גרסאות PTX (compute_90...) נשמרו לצד ה-SASS. לניתוח עמוק יותר של SASS בודד - בניית CFG (control flow graph) ומיפוי חזרה לשורות המקור - משתמשים ב-nvdisasm על ה-cubin שחילצנו, בדיוק כפי שראינו בשיעור 4.2.
גרסאות הקוד המוטמע - compute_XX מול sm_XX¶
עכשיו לשאלה המרכזית: איך אנחנו קובעים מה בדיוק נכנס ל-fatbin? כאן נכנסת ההבחנה שראינו לראשונה בשיעור 2.5, בין הארכיטקטורה הווירטואלית לארכיטקטורה הפיזית:
compute_XX= ארכיטקטורה וירטואלית = PTX, ומספרר לפי compute capability. למשלcompute_90הוא ה-PTX של Hopper (compute capability 9.0).sm_XX= ארכיטקטורה פיזית = SASS, ומספרר לפי גרסת ה-SM. למשלsm_90הוא ה-SASS של ה-H100.
הכלל: PTX תמיד מהודר מארכיטקטורה וירטואלית (compute_XX); SASS תמיד מהודר לארכיטקטורה פיזית (sm_XX). מכאן, ה-PTX הוא נייד קדימה (forward compatible) - PTX של compute_90 יוכל לרוץ על GPU עתידי חדש יותר דרך JIT - ואילו SASS של sm_90 קשור לחומרת Hopper בלבד.
השליטה היא דרך שתי אופציות: --gpu-architecture (קיצור -arch) שקובעת את הארכיטקטורה הווירטואלית לקלט, ו---gpu-code (קיצור -code) שקובעת אילו מטרות לייצר. הצורה המפורשת והחד-משמעית ביותר היא -gencode, שמאגדת את השתיים לזוג אחד ומאפשרת לחזור עליה כמה פעמים:
הטבלה שחשוב לשנן - מה בדיוק נכנס ל-fatbin בכל צורה:
| הדגל | SASS מוטמע | PTX מוטמע | הערה |
|---|---|---|---|
-gencode arch=compute_90,code=sm_90 |
sm_90 |
אין | רק SASS - קטן, מהיר לטעון, אך לא נייד קדימה |
-gencode arch=compute_90,code=compute_90 |
אין | compute_90 |
רק PTX - נייד קדימה, אבל דורש JIT בכל טעינה |
-gencode arch=compute_90,code=\"sm_90,compute_90\" |
sm_90 |
compute_90 |
שניהם - הטוב משני העולמות |
-arch=sm_90 (קיצור) |
sm_90 |
compute_90 |
שקול לשורה שמעליו - מטמיע גם SASS וגם PTX |
-arch=compute_90 |
אין | compute_90 |
PTX בלבד |
נקודה שמפילה רבים (gotcha). קל לחשוב שההבחנה היא פשוט "compute_ זה PTX, sm_ זה SASS", ולסיים בזה - אבל זה לא מדויק בשני מובנים. ראשית, הקיצור -arch=sm_90 אינו שקול ל-code=sm_90 בלבד: הוא מטמיע גם SASS של sm_90 וגם PTX של compute_90 (זו הסיבה שתוכנית שקומפלה עם -arch=sm_90 בלבד עדיין מסוגלת לרוץ על GPU עתידי דרך JIT - יש בה PTX). שנית, גם כשהטמענו code=compute_90 (PTX בלבד, בלי SASS), ה-SASS עדיין ייווצר - פשוט לא על ידי ptxas בזמן ה-build אלא על ידי ה-hardware driver בזמן הריצה, כ-JIT. כלומר PTX לעולם אינו רץ ישירות; מישהו תמיד הופך אותו ל-SASS, השאלה היחידה היא מתי - בזמן build או בזמן טעינה.
פקודה רב-ארכיטקטונית אמיתית. נניח שאנחנו מפיצים תוכנית שצריכה לרוץ על Volta (V100), Ampere (A100) ו-Hopper (H100), וגם על חומרה עתידית שעדיין לא יצאה:
nvcc -gencode arch=compute_70,code=sm_70 \
-gencode arch=compute_80,code=sm_80 \
-gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_90,code=compute_90 \
mm.cu -o mm
מה נכנס ל-fatbin? שלושה cubins של SASS (עבור sm_70, sm_80, sm_90) ועותק PTX אחד (compute_90). ה-cuobjdump -lelf mm יראה שלושה elf, וה-cuobjdump -lptx mm יראה PTX אחד. בזמן ריצה, ה-driver בוחר את ה-cubin המתאים בדיוק לכרטיס: על H100 הוא טוען את sm_90 ישירות (אין JIT); על כרטיס עתידי מדור שאחרי Hopper, שאין לו cubin תואם, הוא נופל אחורה ל-PTX של compute_90 ומבצע JIT אליו. זו בדיוק אסטרטגיית ה"תאימות קדימה" - כוללים SASS למה שמכירים היום, ו-PTX אחד (הגבוה ביותר) כרשת ביטחון למחר.
הערה על סיומות Hopper/Blackwell. כפי שראינו בשיעור 2.5, ל-Hopper יש גם sm_90a (הסיומת a, "architecture-specific") שחושפת פיצ'רים כמו wgmma ו-TMA שאינם ניידים אפילו בתוך אותה גרסת major. אם ה-kernel שלכם משתמש ב-Tensor Core של Hopper דרך אותן הוראות, תצטרכו לקמפל דווקא עם -arch=sm_90a ולא sm_90.
הcompilation מותנה לפי ארכיטקטורה - CUDA_ARCH¶
ראינו ש-nvcc יכול לייצר SASS לכמה ארכיטקטורות מאותו מקור. אבל מה אם אנחנו רוצים קוד מקור שונה לכל ארכיטקטורה - למשל, להשתמש בהוראה שקיימת רק ב-Ampere ומעלה? כאן נכנס המאקרו המוגדר-מראש __CUDA_ARCH__.
שני כללים קובעים את התנהגותו:
- הוא מוגדר רק במעברי הcompilation של הdevice (device passes), ולא במעבר הhost. בקוד הhost (זה שהולך ל-
gcc),__CUDA_ARCH__אינו מוגדר כלל. לכן#ifdef __CUDA_ARCH__הוא הדרך המדויקת להבדיל בין "אני מתקמפל עכשיו לקוד GPU" ל"אני מתקמפל לקוד CPU" בתוך פונקציה שמסומנת__host__ __device__. - ערכו הוא ה-compute capability כפול 100 (major) ועוד 10 (minor). עבור
compute_90הערך הוא900, עבורcompute_80הוא800, עבורcompute_75הוא750. כך אפשר להשוות:#if __CUDA_ARCH__ >= 800.
הנקודה החשובה: כשמקמפלים רב-ארכיטקטונית (רשימת -gencode), nvcc מריץ את cicc מחדש עבור כל ארכיטקטורה וירטואלית, וב-כל מעבר כזה __CUDA_ARCH__ מקבל את הערך המתאים לו. כך אותו קובץ מקור מייצר SASS שונה לכל sm_XX:
__global__ void reduce_step(int* out) {
#if __CUDA_ARCH__ >= 800
// Ampere/Hopper path: hardware reduction instruction at warp level
out[threadIdx.x] = __reduce_add_sync(0xffffffff, threadIdx.x);
#elif __CUDA_ARCH__ >= 700
// Volta/Turing path: classic reduction with shuffle
int v = threadIdx.x;
for (int o = 16; o > 0; o >>= 1) v += __shfl_down_sync(0xffffffff, v, o);
out[threadIdx.x] = v;
#else
out[threadIdx.x] = -1; // unsupported architecture
#endif
}
כשנקמפל את זה עם -gencode arch=compute_70,code=sm_70 -gencode arch=compute_90,code=sm_90, ה-cubin של sm_70 יכיל את מסלול ה-shuffle, וה-cubin של sm_90 יכיל את הוראת REDUX החומרתית - ונוכל לאמת זאת בעינינו עם cuobjdump -sass (זה בדיוק תרגיל 2).
מלכוד נפוץ. מכיוון ש-__CUDA_ARCH__ אינו מוגדר במעבר הhost, אסור להשתמש בו כדי לגדר החלטות שמשפיעות על חתימת פונקציה או על טיפוסים שרואה גם הhost - זה יוצר חוסר-התאמה בין מה שהhost מקמפל למה שהdevice מקמפל, ומוביל לבאגים מסתוריים בקישור או בזמן ריצה. השתמשו בו רק לגידור גוף של קוד device.
הcompilation נפרד וקישור device - separate compilation & device linking¶
עד עכשיו הנחנו שהכל בקובץ .cu אחד. אבל פרויקטים אמיתיים מפוצלים לקבצים רבים, ולעיתים פונקציית __device__ שמוגדרת בקובץ אחד נקראת מ-kernel שבקובץ אחר. כברירת מחדל nvcc עובד ב-whole program compilation (compilation תוכנית-שלמה): כל .cu מהודר בנפרד לגמרי, וקוד הdevice שלו חייב להיות עצמאי. אם ננסה לקרוא __device__ float square(float) שמוגדרת בקובץ אחר, נקבל שגיאת קישור:
כדי לפתור זאת צריך relocatable device code (RDC) - קוד device ניתן-להזזה, שמשאיר את הסמלים פתוחים לקישור מאוחר יותר, בדיוק כמו קבצי .o רגילים בhost. שני שלבים:
// main.cu
#include <cstdio>
extern __device__ float square(float); // declaration, the definition is in another file
__global__ void apply(float* d, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) d[i] = square(d[i]); // cross-file call
}
הcompilation בשלושה שלבים - compilation נפרד לכל קובץ עם -dc, ואז קישור:
nvcc -arch=sm_90 -dc square.cu -o square.o # -dc = --device-c = compilation for RDC
nvcc -arch=sm_90 -dc main.cu -o main.o
nvcc -arch=sm_90 square.o main.o -o app # link step: nvcc runs nvlink automatically
הדגל -dc (קיצור ל---device-c, שקול ל--rdc=true -c) מקמפל כל קובץ לקוד device ניתן-להזזה במקום לקוד תוכנית-שלמה. שלב הקישור האחרון מפעיל את nvlink (הכלי החמישי מהטבלה בסעיף הראשון), שמקשר את סמלי הdevice בין הקבצים - בדיוק כמו ש-ld מקשר סמלי host. אם רוצים לבצע רק את שלב קישור-הdevice במפורש, יש -dlink:
nvcc -arch=sm_90 -dlink square.o main.o -o dev_linked.o # device link only
g++ square.o main.o dev_linked.o -o app -lcudart # then a regular host link
Separate compilation with RDC:
square.cu ─(-dc)─► square.o ┐
├─► nvlink ─► device link ─► executable
main.cu ─(-dc)─► main.o ┘ (links square between the files)
מחיר ה-RDC. קוד ניתן-להזזה עולה קצת בביצועים: ptxas לא יכול לבצע inlining חוצה-קבצים ואופטימיזציות תוכנית-שלמה מסוימות כשהסמלים נשארים פתוחים. לכן RDC מפעילים רק כשבאמת צריך קישור __device__ חוצה-קבצים (או פיצ'רים שדורשים אותו, כמו dynamic parallelism). כשהכל בקובץ אחד, ה-whole program הרגיל מהיר יותר.
הcompilation בזמן ריצה - nvrtc¶
כל מה שראינו עד כה הוא compilation מוקדם (ahead-of-time, AOT): nvcc רץ כתהליך נפרד בזמן ה-build, לפני שהתוכנית בכלל עלתה. אבל מה אם ה-kernel עצמו נוצר בזמן ריצה - למשל framework שמייצר קוד CUDA C++ דינמית לפי הקלט של המשתמש, או מנוע שבונה kernels שונים לפי גדלים שמתגלים רק בזמן הריצה? להריץ nvcc כתת-תהליך בכל פעם יהיה יקר ומסורבל. כאן נכנסת ספריית הcompilation בזמן-ריצה - nvrtc (NVIDIA Runtime Compiler).
ההגדרה: nvrtc היא ספרייה שמהדרת CUDA C++ ל-PTX בתוך התהליך שלכם (in-process), בלי צורך להריץ nvcc בתהליך נפרד. ההבדל מ-nvcc הוא בדיוק בציר AOT מול JIT:
| היבט | nvcc |
nvrtc |
|---|---|---|
| מתי | ahead-of-time, בזמן build | runtime, בזמן ריצה |
| איך | תהליך נפרד | קריאת ספרייה בתוך התהליך |
| פלט | PTX ו/או SASS (fatbin) | PTX בלבד |
| מי הופך ל-SASS | ptxas בזמן build |
ה-hardware driver, ב-JIT |
| קוד פתוח? | חלקי | לא, סגור |
שתי נקודות קריטיות. ראשית, nvrtc עוצרת ב-PTX - היא אינה מייצרת SASS. את השלב השני, PTX -> SASS, מבצע ה-hardware driver של NVIDIA בזמן ריצה, ב-JIT - וזהו אותו מנגנון JIT בדיוק שראינו קודם עבור fatbin שמכיל רק PTX. יש כאן, אם כן, שני compilers נפרדים בשרשרת: (1) nvrtc מהדרת CUDA C++ -> PTX; (2) ה-driver מהדר PTX -> SASS. אל תבלבלו ביניהם. שנית, nvrtc היא קוד סגור (closed source).
CUDA C++ source ──nvrtc──► PTX ──driver JIT──► SASS ──► runs on the SM
(string in memory) (in-process) (at load time)
הנה תוכנית מינימלית מלאה: מחרוזת מקור, nvrtc מהדרת ל-PTX, וה-CUDA Driver API (הinterface ברמה הנמוכה עם התחילית cu, שנעמיק בו בפרק 5) טוענת את ה-PTX (מה שמפעיל את ה-JIT של ה-driver), משיגה את הפונקציה ומשיקה אותה:
#include <nvrtc.h>
#include <cuda.h>
#include <vector>
#include <cstdio>
#include <cstdlib>
#define NVRTC_CHECK(x) \
do { \
nvrtcResult r_ = (x); \
if (r_ != NVRTC_SUCCESS) { \
fprintf(stderr, "NVRTC error %s:%d -> %s\n", __FILE__, __LINE__, \
nvrtcGetErrorString(r_)); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define CU_CHECK(x) \
do { \
CUresult r_ = (x); \
if (r_ != CUDA_SUCCESS) { \
const char* msg_; cuGetErrorName(r_, &msg_); \
fprintf(stderr, "CU error %s:%d -> %s\n", __FILE__, __LINE__, msg_);\
exit(EXIT_FAILURE); \
} \
} while (0)
// the kernel as a source string - created here dynamically at runtime
const char* src = R"cuda(
extern "C" __global__
void saxpy(float a, const float* x, const float* y, float* out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) out[i] = a * x[i] + y[i];
}
)cuda";
int main() {
// ---- Step 1: nvrtc compiles CUDA C++ -> PTX, in-process ----
nvrtcProgram prog;
NVRTC_CHECK(nvrtcCreateProgram(&prog, src, "saxpy.cu", 0, nullptr, nullptr));
const char* opts[] = { "--gpu-architecture=compute_90" }; // PTX target: Hopper
nvrtcResult comp = nvrtcCompileProgram(prog, 1, opts);
size_t logSize; NVRTC_CHECK(nvrtcGetProgramLogSize(prog, &logSize));
if (logSize > 1) { // always print the compilation log
std::vector<char> log(logSize);
NVRTC_CHECK(nvrtcGetProgramLog(prog, log.data()));
fprintf(stderr, "%s\n", log.data());
}
if (comp != NVRTC_SUCCESS) return EXIT_FAILURE;
size_t ptxSize; NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptxSize));
std::vector<char> ptx(ptxSize);
NVRTC_CHECK(nvrtcGetPTX(prog, ptx.data())); // the PTX as a string
NVRTC_CHECK(nvrtcDestroyProgram(&prog));
// ---- Step 2: Driver API loads the PTX (this is where the driver does JIT to SASS) ----
CU_CHECK(cuInit(0));
CUdevice dev; CU_CHECK(cuDeviceGet(&dev, 0));
CUcontext ctx; CU_CHECK(cuCtxCreate(&ctx, 0, dev));
CUmodule mod; CU_CHECK(cuModuleLoadDataEx(&mod, ptx.data(), 0, nullptr, nullptr));
CUfunction fn; CU_CHECK(cuModuleGetFunction(&fn, mod, "saxpy"));
// ---- Step 3: allocation, copy, launch with cuLaunchKernel ----
const int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
std::vector<float> hX(n, 1.0f), hY(n, 2.0f), hOut(n);
CUdeviceptr dX, dY, dOut;
CU_CHECK(cuMemAlloc(&dX, bytes));
CU_CHECK(cuMemAlloc(&dY, bytes));
CU_CHECK(cuMemAlloc(&dOut, bytes));
CU_CHECK(cuMemcpyHtoD(dX, hX.data(), bytes));
CU_CHECK(cuMemcpyHtoD(dY, hY.data(), bytes));
float a = 3.0f;
void* args[] = { &a, &dX, &dY, &dOut, (void*)&n }; // arguments as an array of pointers
int threads = 256, blocks = (n + threads - 1) / threads;
CU_CHECK(cuLaunchKernel(fn, blocks, 1, 1, threads, 1, 1,
0, nullptr, args, nullptr));
CU_CHECK(cuCtxSynchronize());
CU_CHECK(cuMemcpyDtoH(hOut.data(), dOut, bytes));
printf("out[0] = %.1f (expected 5.0)\n", hOut[0]); // 3*1 + 2 = 5
cuMemFree(dX); cuMemFree(dY); cuMemFree(dOut);
cuModuleUnload(mod); cuCtxDestroy(ctx);
return 0;
}
הcompilation וקישור התוכנית עצמה (שימו לב שהיא נקשרת מול libnvrtc ו-libcuda, לא libcudart):
שימו לב לזרימה: אנחנו כותבים את ה-kernel כמחרוזת (שיכולה, כמובן, להיווצר דינמית - להשחיל בה קבועים, גדלים, טיפוסים), מהדרים אותה ל-PTX בזמן ריצה דרך nvrtc, ומעולם לא הרצנו nvcc. שאר הקוד הוא ה-Driver API - cuModuleLoadDataEx הוא הרגע שבו ה-hardware driver מבצע JIT מ-PTX ל-SASS עבור ה-GPU הספציפי שלנו.
שימושים אמיתיים. למה שמישהו יטרח בזה במקום nvcc? כשה-kernel אינו ידוע בזמן ה-build. ספריות כמו CuPy, מנועי inference, ו-JIT של frameworks לרשתות עצביות מייצרות קוד CUDA C++ דינמית - "התמזג (fuse) את שלוש הפעולות האלה ל-kernel אחד עם הקבועים האלה" - ומהדרות אותו on the fly עם nvrtc. זה מאפשר specialization: הטמעת גדלים וקבועים כ-compile-time constants ישירות ב-kernel, מה שנותן ל-ptxas (בשלב ה-JIT) לפרוס לולאות ולייעל בצורה שאי אפשר עם ערכים שמגיעים כארגומנטים בזמן ריצה. הערה: מגרסאות CUDA מודרניות nvrtc יודעת גם לייצר cubin (SASS) ישירות דרך nvrtcGetCUBIN עם --gpu-architecture=sm_90, אבל המודל המנטלי הקלאסי - nvrtc -> PTX, ואז driver JIT -> SASS - הוא הבסיס להבנת המערכת.
סיכום¶
nvccאינו compiler יחיד אלא מנהל-הידור (compiler driver): הוא מפצל את המקור, מוסר את קוד הhost ל-gcc, ומעביר את קוד הdevice דרךcicc(CUDA C++ -> PTX) ו-ptxas(PTX -> SASS), ואורז הכל עםfatbinary.- אל תבלבלו compiler driver (nvcc, gcc - מתזמרים compilation) עם hardware driver (ה-NVIDIA GPU Driver - מדבר עם החומרה); אותה מילה, שני עולמות שונים לגמרי.
- הפלט הוא בינארי host שמציית ל-host ABI ונושא בתוכו fat binary עם PTX ו/או SASS; רואים אותו עם
readelf(סקשן.nv_fatbin) ומפרקים אותו עםcuobjdump -lelf/-lptx/-sass/-ptx. compute_XXמספרר PTX לפי compute capability (ארכיטקטורה וירטואלית);sm_XXמספרר SASS לפי גרסת SM (ארכיטקטורה פיזית); שולטים בהם עם-gencode arch=compute_XX,code=sm_XX.- הקיצור
-arch=sm_90מטמיע גם SASS שלsm_90וגם PTX שלcompute_90; ו-PTX לעולם אינו רץ ישירות - מישהו תמיד הופך אותו ל-SASS, אםptxasבזמן build ואם ה-driver ב-JIT בזמן ריצה. - פקודה רב-ארכיטקטונית עם כמה
-gencodeאורזת cubin לכלsm_XXשמכירים היום ועותק PTX (הגבוה) אחד כרשת ביטחון לחומרה עתידית דרך JIT. __CUDA_ARCH__מוגדר רק במעברי הdevice וערכו הוא ה-compute capability כפול 100 (למשל900ל-Hopper); הוא מאפשר קוד מקור שונה לכל ארכיטקטורה, ואסור להשתמש בו לגידור חתימות שרואה גם הhost.- הcompilation נפרד דורש RDC (
-dc) ו-קישור device (nvlink, נגיש דרך-dlinkאו אוטומטית), כדי לקשר סמלי__device__בין קבצים; במחיר של אובדן אופטימיזציות תוכנית-שלמה. nvrtcמהדרת CUDA C++ ל-PTX בלבד, בתוך התהליך ובזמן ריצה (בליnvcc), וה-driver משלים ל-SASS ב-JIT; שימושית ל-code generation ו-specialization דינמי, וסגורה (closed source).