4.3 nvcc, nvrtc וההידור פתרון
פתרון - nvcc, nvrtc והcompilation¶
הערה: הפלטים למטה הם ממכונת H100 SXM עם CUDA 12.x. הכתובות, מספרי הבתים ופרטי ה-SASS המדויקים ישתנו אצלכם לפי גרסת הכלים והכרטיס (על T4 החליפו sm_90/compute_90 ל-sm_75/compute_75, והערכים ב-__CUDA_ARCH__ יהיו 750). מה שלא ישתנה הוא המבנה: כמה cubins נארזים, מי מוטמע ומי לא, ואיזה מסלול קוד כל ארכיטקטורה לוקחת.
פתרון תרגיל 1 - מה יש בתוך ה-fatbin¶
הקומפילציה הרב-ארכיטקטונית והבדיקה:
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
cuobjdump -lptx mm
readelf -S mm | grep -i nv
הפלט הצפוי:
$ cuobjdump -lelf mm
ELF file 1: mm.1.sm_70.cubin
ELF file 2: mm.2.sm_90.cubin
$ cuobjdump -lptx mm
PTX file 1: mm.1.sm_90.ptx
$ readelf -S mm | grep -i nv
[24] .nv_fatbin PROGBITS ...
[25] __nv_relfatbin PROGBITS ...
שני cubins (SASS ל-sm_70 ול-sm_90) ועותק PTX אחד (compute_90, מסומן ברשימה לפי ה-sm הקרוב). כשמורידים את שורת ה-PTX מהקומפילציה:
nvcc -gencode arch=compute_70,code=sm_70 \
-gencode arch=compute_90,code=sm_90 \
mm.cu -o mm_nptx
cuobjdump -lptx mm_nptx
ולבסוף, הקיצור -arch=sm_90:
$ cuobjdump -lelf mm_arch
ELF file 1: mm_arch.1.sm_90.cubin
$ cuobjdump -lptx mm_arch
PTX file 1: mm_arch.1.sm_90.ptx
למה זה עבד: כל -gencode arch=compute_XX,code=sm_XX מבקש SASS לארכיטקטורה אחת, ולכן קיבלנו שני cubins. השורה code=compute_90 מבקשת עותק PTX, ולכן -lptx הראה אחד; כשהסרנו אותה, לא נותר PTX כלל. ההפתעה בשלב 5: -arch=sm_90 לבדו הטמיע גם SASS של sm_90 וגם PTX של compute_90 - הקיצור שקול ל--gencode arch=compute_90,code="sm_90,compute_90", בדיוק כמו בטבלה בהרצאה. הכל יושב בסקשן ELF בשם .nv_fatbin.
איך להכליל: אסטרטגיית הפצה טובה היא cubin לכל ארכיטקטורה שאתם מכוונים אליה היום (זריזות טעינה, בלי JIT), ועוד עותק PTX אחד של הארכיטקטורה הגבוהה ביותר כרשת ביטחון לחומרה עתידית. בדקו תמיד עם cuobjdump -lelf/-lptx מה באמת נארז - זה מגלה מיד אם שכחתם ארכיטקטורה או הטמעתם PTX מיותר שמנפח את הבינארי.
פתרון תרגיל 2 - התמחות לפי ארכיטקטורה עם CUDA_ARCH¶
// arch.cu
__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
הפלט הצפוי (מקוצר לשורות הרלוונטיות):
code for sm_70
Function : _Z6whoamiPi
...
/*0010*/ MOV R0, 0x2bc ; // 0x2bc = 700 <- the elif path
...
code for sm_90
Function : _Z6whoamiPi
...
/*0010*/ MOV R0, 0x384 ; // 0x384 = 900 <- the if path
...
הרצה על הכרטיס:
למה זה עבד: כשמקמפלים רב-ארכיטקטונית, nvcc מריץ את cicc מחדש לכל ארכיטקטורה וירטואלית, וב-כל מעבר __CUDA_ARCH__ מקבל ערך אחר: 700 במעבר של compute_70 ו-900 במעבר של compute_90. לכן ה-cubin של sm_70 נכנס למסלול ה-elif (טוען 0x2bc = 700), וה-cubin של sm_90 נכנס למסלול ה-if (טוען 0x384 = 900). בזמן ריצה ה-driver בחר את ה-cubin שתואם בדיוק ל-H100, כלומר sm_90, ולכן קיבלנו 900.
איך להכליל: __CUDA_ARCH__ הוא המפתח לכתיבת kernel אחד שמנצל פיצ'רים חדשים כשהם קיימים ונופל אחורה למסלול נייד כשלא (למשל __reduce_add_sync ב-Ampere מול shuffle ב-Volta). תמיד אמתו עם cuobjdump -sass שכל ארכיטקטורה באמת קיבלה את המסלול שהתכוונתם - קל לטעות בתנאי #if ולגלות שכל ה-cubins קיבלו את אותו מסלול.
פתרון תרגיל 3 - compilation מותנה: הdevice מול host¶
// where.cu
#include <cstdio>
__host__ __device__ int where_am_i() {
#ifdef __CUDA_ARCH__
return __CUDA_ARCH__; // device pass
#else
return -1; // host pass
#endif
}
__global__ void probe(int* out) { out[0] = where_am_i(); }
int main() {
printf("host says: %d\n", where_am_i()); // the host
int *d, h;
cudaMalloc(&d, sizeof(int));
probe<<<1, 1>>>(d);
cudaMemcpy(&h, d, sizeof(int), cudaMemcpyDeviceToHost);
printf("device says: %d\n", h); // the device
cudaFree(d);
return 0;
}
הפלט הצפוי:
למה זה עבד: אותה פונקציה where_am_i מתקמפלת פעמיים. במעבר הhost (שהולך ל-gcc) המאקרו __CUDA_ARCH__ אינו מוגדר, ה-#ifdef נכשל, והקוד הוא return -1. במעבר הdevice (עבור sm_90) המאקרו מוגדר לערך 900, ולכן ה-kernel מחזיר 900. זו הדגמה ישירה של הכלל: __CUDA_ARCH__ חי רק בקוד device.
איך להכליל: התבנית #ifdef __CUDA_ARCH__ היא הדרך הסטנדרטית לכתוב פונקציה משותפת __host__ __device__ שרצה על שני הצדדים אבל בוחרת מימוש שונה לכל אחד (למשל sqrtf על ה-GPU מול std::sqrt על הhost). זכרו את המלכוד מההרצאה: אל תגדרו בעזרתו חתימות או טיפוסים שהhost רואה - רק גופי קוד.
פתרון תרגיל 4 - compilation נפרד וקישור device¶
הניסיון בלי RDC נכשל:
הקומפילציה הנכונה עם -dc וקישור 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
הפלט הצפוי:
הבונוס - שלב קישור-הdevice במפורש:
nvcc -arch=sm_90 -dlink square.o app.o -o dev_linked.o
g++ square.o app.o dev_linked.o -o app -L${CUDA_HOME}/lib64 -lcudart
./app
למה זה עבד: כברירת מחדל nvcc מקמפל כל .cu כתוכנית-שלמה סגורה, וההפניה של apply (ב-app.cu) לפונקציה square (ב-square.cu) נשארת לא-פתורה - מכאן ה-Unresolved extern function (השם _Z6squaref הוא ה-mangling של square(float)). הדגל -dc מייצר קוד device ניתן-להזזה (RDC) שמשאיר את הסמל פתוח, ושלב הקישור האחרון מפעיל את nvlink, שמחבר את הסמלים בין הקבצים - בדיוק כמו ld בצד הhost. בבונוס פירקנו את זה: -dlink מייצר אובייקט host שנושא את קוד הdevice המקושר, ואז g++ מבצע את קישור הhost הרגיל מול libcudart.
איך להכליל: הפעילו RDC (-dc/-rdc=true) רק כשאתם באמת צריכים לקרוא __device__ חוצה-קבצים או פיצ'ר שדורש זאת (כמו dynamic parallelism). יש לו מחיר: ptxas מאבד inlining ואופטימיזציות תוכנית-שלמה כשהסמלים פתוחים. אם אפשר לרכז את קוד הdevice בקובץ אחד, ה-whole program הרגיל יהיה מהיר יותר.
פתרון תרגיל 5 - compilation בזמן ריצה עם nvrtc¶
// nvrtc_saxpy.cpp
#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 %s:%d -> %s\n",__FILE__,__LINE__, \
nvrtcGetErrorString(r_)); exit(1);} } while(0)
#define CU_CHECK(x) do { CUresult r_=(x); if(r_!=CUDA_SUCCESS){ \
const char* m_; cuGetErrorName(r_,&m_); \
fprintf(stderr,"CU %s:%d -> %s\n",__FILE__,__LINE__,m_); exit(1);} } while(0)
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() {
// 1. nvrtc: CUDA C++ -> PTX (at runtime, in-process)
nvrtcProgram prog;
NVRTC_CHECK(nvrtcCreateProgram(&prog, src, "saxpy.cu", 0, nullptr, nullptr));
const char* opts[] = { "--gpu-architecture=compute_90" };
nvrtcResult comp = nvrtcCompileProgram(prog, 1, opts);
size_t logSize; NVRTC_CHECK(nvrtcGetProgramLogSize(prog, &logSize));
if (logSize > 1) { std::vector<char> log(logSize);
NVRTC_CHECK(nvrtcGetProgramLog(prog, log.data())); fprintf(stderr, "%s\n", log.data()); }
if (comp != NVRTC_SUCCESS) return 1;
size_t ptxSize; NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptxSize));
std::vector<char> ptx(ptxSize);
NVRTC_CHECK(nvrtcGetPTX(prog, ptx.data()));
NVRTC_CHECK(nvrtcDestroyProgram(&prog));
printf("--- PTX (first lines) ---\n%.200s\n-------------------------\n", ptx.data());
// 2. Driver API: loading 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"));
// 3. allocation, copy, launch
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 };
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]);
cuMemFree(dX); cuMemFree(dY); cuMemFree(dOut);
cuModuleUnload(mod); cuCtxDestroy(ctx);
return 0;
}
הפלט הצפוי:
--- PTX (first lines) ---
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-...
// Cuda compilation tools, release 12.x
//
.version 8.3
.target sm_90
.address_size 64
// .globl saxpy
-------------------------
out[0] = 5.0 (expected 5.0)
למה זה עבד: nvrtc הידרה את מחרוזת המקור ל-PTX בתוך התהליך, בלי להריץ nvcc כלל - שימו לב שה-PTX הוא סתם טקסט, ובראשו .target sm_90 ו-.version (גרסת ה-PTX ISA). ה-extern "C" מנע mangling כך ש-cuModuleGetFunction מצא את הסמל בשם saxpy בדיוק. cuModuleLoadDataEx הוא הרגע שבו ה-hardware driver ביצע JIT מ-PTX ל-SASS עבור ה-H100. קישרנו מול libnvrtc (הcompiler) ו-libcuda (ה-Driver API), ולא מול libcudart, כי כל צד ההרצה כאן הוא Driver API טהור.
איך להכליל: זהו השלד של כל שימוש ב-nvrtc: מחרוזת -> nvrtcCompileProgram -> nvrtcGetPTX -> cuModuleLoadDataEx -> cuLaunchKernel. תמיד משכו ובדקו את יומן הcompilation (nvrtcGetProgramLog) גם בהצלחה - הוא מכיל אזהרות. שימו לב לשני הcompilers בשרשרת: nvrtc ל-PTX, וה-driver ל-SASS; אם משהו נכשל ב-cuModuleLoadDataEx ולא ב-nvrtcCompileProgram, הבעיה היא ב-JIT (למשל compute_XX גבוה מהכרטיס), לא בcompilation המקור.
פתרון תרגיל 6 (בונוס) - specialization דינמי עם nvrtc¶
בונים את מחרוזת המקור דינמית ומטמיעים את N כקבוע compile-time:
char src[1024];
int N = 1 << 20;
snprintf(src, sizeof(src), R"cuda(
extern "C" __global__
void scale(const float* x, float* out) {
const int N = %d; // embedded as a compile-time constant
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
float v = x[i];
#pragma unroll
for (int k = 0; k < 8; k++) v = v * 1.5f; // the loop will unroll - N and 8 are known
out[i] = v;
}
}
)cuda", N);
// ... the same nvrtc + Driver API flow as exercise 5, without passing N as an argument ...
השוואת ה-PTX בין הגרסה עם N/הלולאה כקבועים לגרסה שבה הם ארגומנטים:
Version with embedded constants: the loop disappears entirely (8 multiplies in a row, unrolled),
the bound check is against immediate 0x100000, with no load from an argument.
Version with arguments: the loop remains with a counter, N is loaded from param space,
the bound check is against a register loaded at runtime.
למה זה עבד: כשגודל הווקטור והלולאה ידועים בזמן שבו nvrtc (ואחריו ptxas ב-JIT) מהדרים, הcompiler יכול לפרוס את הלולאה (#pragma unroll על טווח קבוע), לחשב את בדיקת הגבול מול ערך מיידי, ולהיפטר מטעינת פרמטרים - כל אלה בלתי אפשריים כשהערכים מגיעים כארגומנטים ומשתנים בין launch לlaunch. הטמעת הקבוע בזמן יצירת המחרוזת היא בדיוק היתרון שאין לcompilation AOT: אנחנו מייצרים kernel חדש, מותאם לגודל הספציפי, בכל פעם.
איך להכליל: זו הסיבה ש-nvrtc קיים. frameworks כמו CuPy, JAX/XLA, ומנועי inference מייצרים kernels דינמית: הצורה המדויקת של החישוב - הגדלים, הטיפוסים, הפעולות שמתמזגות ל-kernel אחד - מתגלה רק בזמן ריצה, לפי הגרף שהמשתמש בנה. nvcc בזמן build לא יכול היה לקמפל את זה מראש כי ה-kernel פשוט לא היה קיים אז. עם nvrtc, בונים את המקור, מהדרים אותו in-process, וה-driver משלים ל-SASS - וכל התוצאה מהירה כמו kernel שנכתב ידנית לאותו מקרה בדיוק.