5.1 Driver API מול Runtime API פתרון
פתרון - Driver API מול Runtime API¶
להלן פתרונות מלאים לכל התרגילים. המספרים המדויקים בפלט (גרסאות ספריות, נתיבים, גרסת CUDA) ישתנו לפי המכונה, הדרייבר וגרסת ה-Toolkit שלכם - התמקדו במבנה של הפלט ולא בספרות המדויקות. כל הפלטים למטה צולמו על מכונת H100 עם CUDA 12.x, אך התבנית זהה על כל GPU של NVIDIA.
פתרון תרגיל 1 - baseline של ה-Runtime API¶
הcompilation וההרצה:
מרגע ההקצאה ועד הlaunch נדרשו חמש קריאות cuda*: שלוש cudaMalloc, שתי cudaMemcpy, ואז הlaunch <<<>>> (שהיא עצמה סוכר תחבירי מעל קריאת Runtime). ה-context של הdevice נוצר בשקט בקריאה הראשונה שנוגעת ב-GPU - כאן, cudaMalloc הראשון. אין בקוד שום קריאת אתחול מפורשת: זה בדיוק היופי (והמלכוד) של ניהול ה-context המרומז של ה-Runtime.
למה זה עבד: ה-Runtime API יוצר ומאתחל את ה-primary context של הdevice באופן עצל (lazily), בקריאה הראשונה שדורשת GPU. המפתח לעולם אינו רואה את השלב הזה.
איך להכליל: כל תוכנית Runtime "פשוטה" מסתירה בתוכה אתחול context שלם. כשנרד ל-Driver API בתרגיל 3, בדיוק השלב הזה יהפוך למפורש - וזו ההשוואה שכל התרגול בנוי סביבה.
פתרון תרגיל 2 - זיהוי השכבות עם ldd¶
הבינארי בברירת המחדל:
libcudart אינו מופיע, כי ברירת המחדל של nvcc היא --cudart=static - ה-Runtime מוטמע בתוך הבינארי. אבל libcuda.so.1 (הדרייבר) כן מופיע, כי ה-Runtime בסופו של דבר עוטף וקורא ל-Driver, שנשאר תמיד ספרייה דינמית חיצונית.
עכשיו עם קישור Runtime דינמי מפורש:
nvcc --cudart shared -arch=sm_90 runtime_vecadd.cu -o runtime_shared
ldd runtime_shared | grep -E 'cuda|cudart'
libcudart.so.12 => /usr/local/cuda/lib64/libcudart.so.12 (0x00007f...)
libcuda.so.1 => /lib/x86_64-linux-gnu/libcuda.so.1 (0x00007f...)
עכשיו שתי השכבות גלויות: libcudart.so.12 (Runtime, מגיע עם ה-CUDA Toolkit, נמצא תחת /usr/local/cuda/lib64) מעל libcuda.so.1 (Driver, מגיע עם דרייבר NVIDIA, נמצא תחת נתיבי מערכת כמו /lib/x86_64-linux-gnu).
למה זה עבד: ספרייה מקושרת-סטטית מוטמעת בבינארי ולא מופיעה ב-ldd; ספרייה מקושרת-דינמית מופיעה כרשומת תלות שה-loader פותר בזמן ריצה. ברירת המחדל הסטטית של nvcc מסתירה את libcudart, אבל את libcuda.so אי אפשר לקשר סטטית - הוא תמיד דינמי.
איך להכליל: ldd הוא כלי האבחון הראשון לכל בינארי CUDA. הנוכחות של libcuda.so.1 היא הסימן שהדרייבר מותקן; הנוכחות של libcudart.so (כשהיא דינמית) מצביעה על תלות ב-Toolkit. אם libcuda.so.1 חסר או "not found", הדרייבר לא הותקן כראוי.
פתרון תרגיל 3 - אותו kernel דרך ה-Driver API¶
ראשית הcompilation ל-PTX:
תוכנית ה-Driver API המלאה:
// driver_vecadd.cpp
#include <cuda.h>
#include <vector>
#include <cstdio>
#include <cstdlib>
#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)
static std::vector<char> readFile(const char* path) {
FILE* f = fopen(path, "rb");
if (!f) { perror("fopen"); exit(EXIT_FAILURE); }
fseek(f, 0, SEEK_END); long sz = ftell(f); fseek(f, 0, SEEK_SET);
std::vector<char> buf(sz + 1);
fread(buf.data(), 1, sz, f); buf[sz] = '\0'; // PTX must end in null
fclose(f);
return buf;
}
int main() {
const int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
std::vector<float> ha(n, 1.0f), hb(n, 2.0f), hc(n);
CU_CHECK(cuInit(0)); // explicit init - no implicit magic
CUdevice dev; CU_CHECK(cuDeviceGet(&dev, 0));
CUcontext ctx; CU_CHECK(cuCtxCreate(&ctx, 0, dev));
std::vector<char> ptx = readFile("kernel.ptx");
CUmodule mod; CU_CHECK(cuModuleLoadData(&mod, ptx.data())); // manual load + JIT
CUfunction fn; CU_CHECK(cuModuleGetFunction(&fn, mod, "vecAdd"));
CUdeviceptr da, db, dc;
CU_CHECK(cuMemAlloc(&da, bytes));
CU_CHECK(cuMemAlloc(&db, bytes));
CU_CHECK(cuMemAlloc(&dc, bytes));
CU_CHECK(cuMemcpyHtoD(da, ha.data(), bytes));
CU_CHECK(cuMemcpyHtoD(db, hb.data(), bytes));
int block = 256, grid = (n + block - 1) / block;
void* args[] = { &da, &db, &dc, (void*)&n }; // arguments as an array of pointers
CU_CHECK(cuLaunchKernel(fn, grid, 1, 1, block, 1, 1,
0, nullptr, args, nullptr));
CU_CHECK(cuCtxSynchronize());
CU_CHECK(cuMemcpyDtoH(hc.data(), dc, bytes));
printf("c[0] = %.1f c[n-1] = %.1f (expected 3.0)\n", hc[0], hc[n-1]);
cuMemFree(da); cuMemFree(db); cuMemFree(dc);
cuModuleUnload(mod); cuCtxDestroy(ctx);
return 0;
}
והשוואת ה-ldd:
התוכנית נקשרת רק ל-libcuda.so.1, בלי libcudart כלל - כי לא קראנו לשום cuda*, רק ל-cu*. זו תלות מינימלית: דרייבר בלבד, בלי Toolkit.
למה זה עבד: השרשרת המפורשת של ה-Driver API עשתה ידנית בדיוק את מה שה-Runtime עשה בשקט בתרגיל 1 - אתחול (cuInit), יצירת context (cuCtxCreate), טעינת קוד ה-kernel (cuModuleLoadData, שכאן מפעיל את ה-JIT של הדרייבר מ-PTX ל-SASS של sm_90), וlaunch (cuLaunchKernel). התוצאה זהה כי זה אותו kernel על אותה חומרה.
איך להכליל: הדפוס cuInit -> cuDeviceGet -> cuCtxCreate -> cuModuleLoad(Data) -> cuModuleGetFunction -> cuLaunchKernel הוא השלד של כל שימוש ב-Driver API, כולל טעינת PTX שנוצר ב-nvrtc (שיעור 4.3). ברגע שקוד ה-GPU מגיע ממקור חיצוני או דינמי, זו הדרך היחידה לטעון ולהשיק אותו.
פתרון תרגיל 4 - הבטחת התאימות קדימה¶
ניסוח ההבטחה. יישום שקומפל מול גרסה ישנה של ה-Driver API ירוץ על מערכת עם דרייבר (libcuda.so) חדש יותר. הצד הישן הוא האפליקציה; הצד החדש הוא הדרייבר על מכונת היעד. הכיוון ההפוך - אפליקציה שקומפלה מול Toolkit חדש על מכונה עם דרייבר ישן - אינו מובטח, כי האפליקציה עלולה לדרוש סמל שעדיין לא קיים בדרייבר הישן.
שלושת המנגנונים והראיות:
- קישור דינמי ל-
libcuda.so. הבינארי מכיל רק רשומת תלות בשםlibcuda.so.1, וה-loader פותר אותה מול הדרייבר המותקן בזמן ריצה. הראיה:ldd driver_vecaddמתרגיל 3 מראהlibcuda.so.1 => /lib/x86_64-linux-gnu/libcuda.so.1- כתובת שנפתרת בזמן ריצה, לא קוד מוטמע. - תאימות ABI קדימה. כל סמל שהיה קיים ב-
libcuda.soישן קיים גם בחדש, עם אותה סמנטיקה, ולכן בינארי ישן מוצא את כל ה-cu*שהוא צריך. זו הבטחה של NVIDIA לגבי יציבות ה-ABI של הדרייבר. - JIT של PTX ל-SASS. אם הבינארי נושא PTX, גם ה-GPU יכול להיות חדש יותר: הדרייבר יבצע JIT מה-PTX ל-SASS של החומרה החדשה. הראיה:
יש PTX מוטמע (compute_90) שיכול לעבור JIT קדימה על חומרה מדור עתידי.
גרסת דרייבר מול גרסת Toolkit. הריצו:
| NVIDIA-SMI 550.xx.xx Driver Version: 550.xx.xx CUDA Version: 12.4 |
Cuda compilation tools, release 12.2, V12.2.140
ה-CUDA Version ש-nvidia-smi מציג הוא הגרסה המקסימלית של ה-Runtime API שהדרייבר המותקן תומך בה - כלומר מאפיין של הדרייבר, לא של ה-Toolkit. ה-nvcc --version מציג את גרסת ה-Toolkit שאיתה מקמפלים. תאימות קדימה אומרת שאפשר ש-nvcc (Toolkit) יהיה ישן מהדרייבר - כאן 12.2 מול 12.4 - וזה בסדר. אם היה הפוך (Toolkit 12.4 על דרייבר שתומך רק ב-12.2), היינו עלולים לקבל את שגיאת ה-insufficient driver version.
למה זה עבד: שלושת המנגנונים יחד מנתקים את הבינארי מגרסה קונקרטית של הדרייבר ושל החומרה. הבינארי מתאר "מה הוא צריך" (סמלים ב-libcuda.so, PTX) ולא "איפה זה נמצא" או "לאיזו חומרה", וההכרעות האלה נדחות לזמן ריצה.
איך להכליל: זו הסיבה שכדאי תמיד לכלול PTX ב-fatbin (רשת ביטחון לחומרה עתידית), ולקמפל מול ה-Toolkit הישן ביותר שאתם מתחייבים אליו. בפריסה, "דרייבר חדש מריץ בינארי ישן" הוא הכיוון הבטוח.
פתרון תרגיל 5 - מתי צריך את ה-Driver API¶
- מנוע inference שמייצר קוד בזמן ריצה - חובה Driver API. קוד ה-GPU אינו ידוע בזמן build, ולכן חייבים
cuModuleLoadDataכדי לטעון PTX שנוצר בזמן ריצה (בדרך כלל דרך nvrtc, כמו בשיעור 4.3). ה-<<<>>>של ה-Runtime דורש kernel שקומפל מראש - לא רלוונטי כאן. - כפל מטריצות בגודל קבוע - Runtime API מספיק. ה-kernel ידוע בזמן compilation, אין צורך בשליטה מפורשת ב-context. זה בדיוק המקרה הנפוץ שה-Runtime נועד לו; ה-Driver API רק יוסיף מאות שורות boilerplate ללא תועלת.
- כלי לפריסה עם דרייבר בלבד - נוטה חזק ל-Driver API. קישור רק ל-
libcuda.so(שקיים תמיד עם הדרייבר) מסיר את התלות ב-libcudart.soובכל ה-Toolkit, ונותן תלות פריסה מינימלית. זה שיקול תפעולי מובהק לטובת ה-Driver API. - שני contexts מבודדים על אותו GPU - חובה Driver API. ה-Runtime מנהל primary context יחיד ומשותף ואינו נותן ליצור contexts מבודדים;
cuCtxCreateשל ה-Driver API הוא הדרך היחידה לבידוד מפורש כזה.
למה זה עבד: שתי שאלות ההכרעה - "האם קוד ה-GPU ידוע בזמן build?" ו"האם צריך שליטה מפורשת ב-context / תלות מינימלית?" - ממפות כל תרחיש לשכבה הנכונה. תרחישים 1, 3, 4 עונים "לא/כן" ולכן דורשים Driver API; תרחיש 2 הוא Runtime קלאסי.
איך להכליל: ברירת המחדל היא תמיד Runtime API. יורדים ל-Driver API רק כשנתקלים בקיר קונקרטי - קוד דינמי, בידוד context, או גיזום תלות - שה-Runtime לא נותן לטפס עליו.
פתרון תרגיל 6 (בונוס) - אינטרופ בין שתי הinterfaces¶
// interop.cu - compile: nvcc -arch=sm_90 interop.cu -o interop -lcuda
#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>
#include <cstdio>
#include <cstdlib>
#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)
#define RT_CHECK(x) do { cudaError_t e_=(x); if(e_!=cudaSuccess){ \
fprintf(stderr,"RT %s:%d %s\n",__FILE__,__LINE__,cudaGetErrorString(e_)); \
exit(1);} } while(0)
extern "C" __global__ void vecAdd(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];
}
int main() {
const int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
std::vector<float> ha(n, 1.0f), hb(n, 2.0f), hc(n);
CU_CHECK(cuInit(0));
CUdevice dev; CU_CHECK(cuDeviceGet(&dev, 0));
CUcontext primary;
CU_CHECK(cuDevicePrimaryCtxRetain(&primary, dev)); // the same context the Runtime uses
CU_CHECK(cuCtxSetCurrent(primary));
// allocation through the Runtime (cuda*)
float *da, *db, *dc;
RT_CHECK(cudaMalloc(&da, bytes));
RT_CHECK(cudaMalloc(&db, bytes));
RT_CHECK(cudaMalloc(&dc, bytes));
RT_CHECK(cudaMemcpy(da, ha.data(), bytes, cudaMemcpyHostToDevice));
RT_CHECK(cudaMemcpy(db, hb.data(), bytes, cudaMemcpyHostToDevice));
// launch through the Driver (cu*) on the same pointers.
// you cannot simply cast the Runtime's symbol to a CUfunction - it is an opaque handle
// obtained only from cuModuleGetFunction. Get a valid CUfunction from the Runtime's symbol
// via cudaGetFuncBySymbol (cudaFunction_t and CUfunction are the same underlying type).
CUfunction fn;
RT_CHECK(cudaGetFuncBySymbol((cudaFunction_t*)&fn, (const void*)vecAdd));
int block = 256, grid = (n + block - 1) / block;
CUdeviceptr pa=(CUdeviceptr)da, pb=(CUdeviceptr)db, pc=(CUdeviceptr)dc;
void* args[] = { &pa, &pb, &pc, (void*)&n };
CU_CHECK(cuLaunchKernel(fn, grid,1,1, block,1,1, 0, nullptr, args, nullptr));
CU_CHECK(cuCtxSynchronize());
RT_CHECK(cudaMemcpy(hc.data(), dc, bytes, cudaMemcpyDeviceToHost));
printf("interop c[0] = %.1f (expected 3.0)\n", hc[0]);
cudaFree(da); cudaFree(db); cudaFree(dc);
CU_CHECK(cuDevicePrimaryCtxRelease(dev));
return 0;
}
כשמשתמשים ב-cuDevicePrimaryCtxRetain, ה-Runtime (cudaMalloc) וה-Driver (cuLaunchKernel) חולקים את אותו primary context, ולכן המצביעים שהוקצו בצד ה-Runtime תקפים בצד ה-Driver והlaunch מצליחה. אם במקום זאת היינו יוצרים context חדש ונפרד עם cuCtxCreate ומגדירים אותו כ-current, ה-cudaMalloc היה עדיין רץ על ה-primary context (כי ככה ה-Runtime תמיד עובד), בעוד ה-cuLaunchKernel היה רץ על ה-context החדש - שני מרחבי כתובות נפרדים. המצביע מה-Runtime לא היה נראה ב-context של ה-Driver, והתוצאה שגיאה כמו CUDA_ERROR_ILLEGAL_ADDRESS או תוצאה שגויה.
למה זה עבד: ה-context הוא גבול מרחב הכתובות של ה-GPU. הקצאה שייכת ל-context שבו נעשתה, ונראית רק בו. שיתוף ה-primary context בין שתי הinterfaces הוא מה שמאפשר להן לדבר על אותו זיכרון.
איך להכליל: בכל פעם שמערבבים Runtime ו-Driver API באותו thread - למשל שילוב ספרייה שכתובה מול ה-Driver API עם קוד יישום שכתוב מול ה-Runtime - חייבים לוודא ששניהם על ה-primary context (cuDevicePrimaryCtxRetain), אחרת ההקצאות לא ייראו זו לזו. זהו ה-caveat המרכזי של ניהול ה-context שההרצאה הזהירה ממנו.