4.3 nvcc, nvrtc וההידור תרגול

תרגול - nvcc, nvrtc והcompilation¶

בתרגול הזה תפרקו את שרשרת הcompilation של CUDA במו ידיכם. תתחילו בהצצה לתוך ה-fatbin - תקמפלו לשתי ארכיטקטורות ותראו בעיניכם אילו cubins ואילו גרסאות PTX נארזו פנימה. אחר כך תשתמשו ב-__CUDA_ARCH__ כדי שאותו קובץ מקור ייצור SASS שונה לכל ארכיטקטורה, ותאמתו את זה מול ה-SASS עצמו. תפצלו פונקציית __device__ בין שני קבצים ותקשרו אותם עם -dc/-dlink, תגלו מה קורה בלי RDC, ולבסוף תכתבו תוכנית nvrtc מלאה שמהדרת kernel ממחרוזת בזמן ריצה ומשיקה אותו דרך ה-Driver API. עבדו לפי הסדר - כל תרגיל נשען על ההבנה מהקודם. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס (sm_90), אבל כל תרגיל כאן עובד על כל GPU של NVIDIA - רק החליפו את מספרי הארכיטקטורה לאלה של הכרטיס שלכם.

הכנה¶

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

nvidia-smi --query-gpu=name,compute_cap --format=csv
# e.g.:  NVIDIA H100 80GB HBM3, 9.0   -> sm_90 / compute_90
#         Tesla T4, 7.5                 -> sm_75 / compute_75

לאורך התרגול נשתמש ב-kernel פשוט של כפל מטריצות מהיר (או כל kernel אחר). צרו את mm.cu:

#include <cstdio>
__global__ void mm(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float s = 0.0f;
        for (int k = 0; k < N; k++) s += A[row * N + k] * B[k * N + col];
        C[row * N + col] = s;
    }
}
int main() { printf("stub\n"); return 0; }

בכל מקום שמופיע sm_90/compute_90 בתרגילים, החליפו לזוג של הכרטיס שלכם (למשל sm_75/compute_75 ל-T4). את הזוג הנמוך בתרגילים הרב-ארכיטקטוניים (sm_70) אפשר להשאיר - רק דאגו שהוא נמוך או שווה לכרטיס שלכם, כדי ש-ptxas יסכים לקמפל.

תרגיל 1 - מה יש בתוך ה-fatbin¶

קמפלו את mm.cu לשתי ארכיטקטורות בו-זמנית, עם SASS לשתיהן ו-PTX של הגבוהה כרשת ביטחון:

nvcc -gencode arch=compute_70,code=sm_70 \
     -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     mm.cu -o mm

השתמשו ב-cuobjdump -lelf mm כדי לרשום את ה-cubins (SASS) הארוזים, וב-cuobjdump -lptx mm כדי לרשום את גרסאות ה-PTX. כמה cubins יש? כמה PTX? לאילו ארכיטקטורות?
הריצו readelf -S mm | grep -i nv וזהו את הסקשן שנושא את קוד ה-GPU. איזה שם יש לו?
עכשיו קמפלו גרסה בלי ה--gencode של ה-PTX (רק שני ה-SASS), והשוו את פלט cuobjdump -lptx. מה השתנה?
קמפלו גרסה שלישית עם הקיצור -arch=sm_90 בלבד, והריצו עליה גם -lelf וגם -lptx. מה מצאתם, ולמה זה מפתיע?

רמז: -lelf ו--lptx הם "list" - הם רק מונים ומתארים, בלי להדפיס את הקוד עצמו (לזה יש -sass ו--ptx). בשלב 5 שימו לב: הקיצור -arch=sm_90 לבדו מטמיע גם SASS של sm_90 וגם PTX של compute_90 - בדיוק כמו -gencode arch=compute_90,code="sm_90,compute_90". זו הנקודה מהטבלה בהרצאה. הסקשן ב-ELF נקרא .nv_fatbin.

תרגיל 2 - התמחות לפי ארכיטקטורה עם CUDA_ARCH¶

כתבו kernel שכותב ערך שונה לכל ארכיטקטורה, כדי ליצור מבדיל שקל לזהות ב-SASS:

__global__ void whoami(int* out) {
#if __CUDA_ARCH__ >= 900
    out[0] = 900;
#elif __CUDA_ARCH__ >= 700
    out[0] = 700;
#else
    out[0] = 0;
#endif
}

קמפלו לשתי ארכיטקטורות: nvcc -gencode arch=compute_70,code=sm_70 -gencode arch=compute_90,code=sm_90 -c arch.cu -o arch.o.
הריצו cuobjdump -sass arch.o והסתכלו על ה-SASS של שני ה-cubins. חפשו את ההוראה שטוענת את הקבוע לפני הכתיבה ל-out[0]. איזה ערך מיידי (immediate) מופיע בכל cubin?
אמתו: 900 בהקסה הוא 0x384, ו-700 הוא 0x2bc. ודאו שכל cubin טוען את הערך הנכון לו.
הריצו את ה-kernel על הכרטיס שלכם (השיקו whoami<<<1,1>>> והעתיקו את out[0] חזרה). איזה ערך קיבלתם, ולמה דווקא אותו?

רמז: cuobjdump -sass מציג את ה-SASS מקובץ ל-cubins, כל אחד עם כותרת שמציינת את ה-sm_XX שלו. חפשו הוראת MOV (או IMAD.MOV/ULDC) עם הקבוע. לגבי שלב 5: ה-driver בוחר את ה-cubin שתואם בדיוק לכרטיס שלכם - על H100 זה ה-sm_90 (ותקבלו 900), על T4 זה ה-sm_75 (שנופל ל-elif, ותקבלו 700). המרה: printf("%d\n", 0x384) נותן 900.

תרגיל 3 - compilation מותנה: הdevice מול host¶

כתבו פונקציה אחת עם __host__ __device__ שמתנהגת שונה בשני העולמות:

#include <cstdio>
__host__ __device__ int where_am_i() {
#ifdef __CUDA_ARCH__
    return __CUDA_ARCH__;   // compiles to GPU code
#else
    return -1;              // compiles to CPU code
#endif
}
__global__ void probe(int* out) { out[0] = where_am_i(); }

קראו ל-where_am_i() גם מהhost (ב-main, והדפיסו), וגם דרך ה-kernel probe (העתיקו את out[0] חזרה והדפיסו).
קמפלו עם -arch=sm_90 והריצו. איזה ערך הדפיס הhost ואיזה ה-kernel?
הסבירו: למה אותה פונקציה בדיוק החזירה שני ערכים שונים?

רמז: במעבר הhost __CUDA_ARCH__ אינו מוגדר, ולכן ה-#ifdef נכשל והמסלול הוא return -1. במעבר הdevice הוא מוגדר לערך של הארכיטקטורה (900 ל-sm_90), ולכן ה-kernel מחזיר 900. זו הדגמה של הכלל מההרצאה: __CUDA_ARCH__ מוגדר רק בקוד device. הhost ידפיס -1, ה-kernel ידפיס 900.

תרגיל 4 - compilation נפרד וקישור device¶

פצלו פונקציית __device__ לקובץ נפרד. square.cu:

__device__ float square(float x) { return x * x; }

ו-app.cu:

#include <cstdio>
extern __device__ float square(float);
__global__ void apply(float* d, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) d[i] = square(d[i]);
}
int main() {
    int n = 8; float h[8]; for (int i=0;i<n;i++) h[i]=(float)i;
    float* d; cudaMalloc(&d, n*sizeof(float));
    cudaMemcpy(d, h, n*sizeof(float), cudaMemcpyHostToDevice);
    apply<<<1, n>>>(d, n);
    cudaMemcpy(h, d, n*sizeof(float), cudaMemcpyDeviceToHost);
    for (int i=0;i<n;i++) printf("%.0f ", h[i]); printf("\n");
    cudaFree(d); return 0;
}

קודם נסו לקמפל בלי RDC, בדרך הרגילה, וראו את השגיאה:
```
nvcc -arch=sm_90 square.cu app.cu -o app
```
מה השגיאה, ולמה היא קורה?

עכשיו קמפלו עם compilation נפרד וקישור device:

nvcc -arch=sm_90 -dc square.cu -o square.o
nvcc -arch=sm_90 -dc app.cu    -o app.o
nvcc -arch=sm_90 square.o app.o -o app

הריצו את app. מה הפלט הצפוי?

(בונוס) בצעו את שלב הקישור האחרון בשני שלבים נפרדים - -dlink ואז קישור host עם g++ - כדי לראות את שלב קישור-הdevice במפורש.

רמז: בלי -dc, כל קובץ מהודר כתוכנית-שלמה עצמאית, וההפניה של apply ל-square שבקובץ אחר נשארת לא-פתורה. השגיאה תהיה בסגנון ptxas fatal : Unresolved extern function או nvlink error. הדגל -dc (שווה ל--rdc=true -c) מייצר קוד device ניתן-להזזה, ושלב הקישור האחרון מפעיל את nvlink. הפלט הצפוי הוא הריבועים: 0 1 4 9 16 25 36 49.

תרגיל 5 - compilation בזמן ריצה עם nvrtc¶

כתבו תוכנית שמהדרת kernel ממחרוזת בזמן ריצה עם nvrtc, ומשיקה אותו דרך ה-Driver API. השתמשו ב-kernel של saxpy (out = a*x + y).
השלבים בקוד: nvrtcCreateProgram -> nvrtcCompileProgram (עם --gpu-architecture=compute_90) -> תמיד משכו ומדפיסים את יומן הcompilation -> nvrtcGetPTX. אחר כך Driver API: cuInit -> cuCtxCreate -> cuModuleLoadDataEx (כאן ה-driver עושה JIT) -> cuModuleGetFunction -> cuLaunchKernel.
קמפלו וקשרו מול הספריות הנכונות:
```
nvcc nvrtc_saxpy.cpp -o nvrtc_saxpy -lnvrtc -lcuda
```
הריצו ואמתו את התוצאה (a=3, x=1, y=2 -> out=5).
הוסיפו הדפסה של ה-PTX שקיבלתם מ-nvrtc (הוא סתם מחרוזת), וזהו בו את שם ה-kernel ואת הגרסה (.target sm_90 / .version).

רמז: השתמשו בקוד המלא מההרצאה כנקודת פתיחה. שימו לב לשלושה דברים: (א) ה-kernel חייב להיות extern "C" כדי שהשם saxpy לא יעבור mangling וש-cuModuleGetFunction ימצא אותו; (ב) קשרו מול -lnvrtc -lcuda (לא libcudart - זו תוכנית Driver API); (ג) הארגומנטים ל-cuLaunchKernel עוברים כמערך של מצביעים ל-void, void* args[] = {&a, &dX, ...}. אם ה-JIT נכשל, בדקו שגרסת ה-compute_XX שביקשתם אינה גבוהה מהכרטיס.

תרגיל 6 (בונוס) - specialization דינמי עם nvrtc¶

הרחיבו את תוכנית ה-nvrtc כך שתייצרו את מחרוזת ה-kernel דינמית: הטמיעו את גודל הווקטור n כקבוע compile-time בתוך המקור (למשל בנו את המחרוזת עם snprintf והכניסו const int N = 1048576;), במקום להעביר את n כארגומנטים.
קמפלו את שתי הגרסאות (עם n כארגומנט מול n מוטמע כקבוע) והשוו את ה-PTX שמפיק nvrtc. האם הלולאה (אם הוספתם כזו) נפרסת אחרת?
הסבירו: למה הטמעת קבועים ב-compile-time מאפשרת ל-ptxas (בשלב ה-JIT) לייצר קוד יעיל יותר מאשר ערכים שמגיעים כארגומנטים?
(מחשבה) איזה סוג של framework היה מרוויח מזה בעולם האמיתי, ולמה nvcc רגיל לא היה מתאים לתרחיש הזה?

רמז: כשגודל ידוע ב-compile-time, ה-JIT יכול לפרוס לולאות, לחשב היסטים מראש, ולהיפטר מבדיקות גבול - כל מה שאי אפשר כשהערך מגיע כארגומנט בזמן ריצה. זו בדיוק הסיבה ש-frameworks כמו CuPy ומנועי inference מייצרים kernels דינמית: ה-kernel המדויק (הגדלים, הטיפוסים, הפעולות המתמזגות) ידוע רק בזמן ריצה, ולכן nvcc בזמן build לא יכול היה לקמפל אותו מראש.