לדלג לתוכן

5.1 Driver API מול Runtime API הרצאה

בשיעור 0.4 ציירנו את מפת שרשרת התוכנה כולה, ובתחתית נתיב ההרצה ראינו שלוש שכבות שמפרידות בין היישום שלכם לבין הסיליקון: ה-Runtime API, ה-Driver API, ו-nvidia.ko. בשיעור 4.3, כשבנינו kernel (קרנל - פונקציית ה-GPU) בזמן ריצה עם nvrtc, כבר טעמנו את ה-Driver API מקרוב - קראנו ל-cuModuleLoadDataEx ול-cuLaunchKernel בלי להסביר לעומק מדוע. השיעור הזה פותח את פרק 5, פרק מחסנית ה-host (host - צד ה-CPU), והוא מוקדש כולו לשתי הinterfaces שהיישום שלכם מדבר דרכם אל ה-GPU (מעבד גרפי / כרטיס מסך): ה-CUDA Runtime API הגבוה והנפוץ, וה-CUDA Driver API הנמוך והמפורש. נבין בדיוק איך הם מסודרים בשכבות, מי עוטף את מי, מהו ההבדל בניהול ה-context ובlaunch ה-kernel, מהי הבטחת התאימות קדימה שמאפשרת לכם לשלוח בינארי היום ולהריץ אותו על דרייבר של מחר, ומתי באמת - אם בכלל - כדאי לרדת ישירות ל-Driver API. את מה שיושב מתחת לשתי הספריות האלה, כלומר את הדרייבר עצמו על שני חלקיו, נפרק בשיעור 5.2.

שתי הinterfaces ושכבות המחסנית - the two host APIs

נתחיל מהתמונה הכוללת. כל בקשה שהיישום שלכם שולח ל-GPU עוברת דרך מדרג (stack) קבוע של שכבות, כל אחת מפשטת (abstracts) את זו שמתחתיה. שתי השכבות העליונות הן interfaces תכנות - APIs - ושתיהן ממומשות כספריות user mode (מרחב המשתמש) שרצות בתוך התהליך שלכם:

  +---------------------------------------------------------+
  |  the application  -  your host code                       |   user mode
  +---------------------------------------------------------+
        |  calls the API. "in order" (NVIDIA's order):
        v
  +---------------------------------------------------------+
  |  CUDA Runtime API   cuda*    libcudart.so   <- high      |   user mode
  |  implicit context management, ergonomic, most common      |
  +---------------------------------------------------------+
        |  the Runtime wraps the Driver
        v
  +---------------------------------------------------------+
  |  CUDA Driver API    cu*      libcuda.so     <- low       |   user mode
  |  explicit context, explicit module loading, fine control  |
  +---------------------------------------------------------+
   -------------------- user / kernel boundary ----------------
        |  (syscall / ioctl)
        v
  +---------------------------------------------------------+
  |  nvidia.ko   kernel module, privileged mode                |   kernel mode
  +---------------------------------------------------------+
        |  touches the hardware directly
        v
  +---------------------------------------------------------+
  |  the GPU hardware (device):  132 SMs on H100, HBM3 80 GiB  |   silicon
  +---------------------------------------------------------+

שתי עובדות יסוד לגבי המבנה הזה. ראשית, ה-Runtime API אינו interface עצמאי היושב ליד ה-Driver API - הוא יושב מעליו ועוטף אותו. כמעט כל קריאת Runtime מתורגמת, בסופו של דבר, לקריאת Driver אחת או יותר. שנית, הביטוי של NVIDIA "הinterfaces לדרייבר הם, לפי הסדר, ה-Runtime API וה-Driver API" מקפל בתוכו את שני הדברים גם יחד: גם את סדר ההעדפה (התחילו ב-Runtime, רדו ל-Driver רק כשצריך) וגם את סדר השכבות (ה-Runtime מעל ה-Driver). ה-Driver API הוא "פתח המילוט" ברמה הנמוכה, לא ברירת המחדל.

ההבחנה הכי מהירה בין השניים היא הקידומת של שמות הפונקציות, וכדאי לחקוק אותה עכשיו:

מאפיין CUDA Runtime API CUDA Driver API
קידומת פונקציות cuda* (למשל cudaMalloc) cu* (למשל cuMemAlloc)
ספרייה ממַמֶשת (Linux) libcudart.so libcuda.so
מגיעה עם CUDA Toolkit הדרייבר של NVIDIA
רמת הפשטה גבוהה, ארגונומית נמוכה, שליטה עדינה
ניהול context מרומז (implicit) מפורש (explicit)
טעינת קוד ה-kernel אוטומטית מה-fatbin ידנית (cuModuleLoad)
קישור (linking) סטטי או דינמי דינמי בלבד
קוד פתוח סגור סגור
מי כותב מולו כמעט כולם מעטים מאוד

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

ה-Runtime API - הinterface הגבוה

ה-CUDA Runtime API הוא הinterface שרוב מוחלט של תוכניות CUDA, וכמעט כל הספריות (cuBLAS, cuDNN, PyTorch), נכתבות מולו. את כל הקוד שכתבנו עד כה בקורס - cudaMalloc, cudaMemcpy, cudaDeviceSynchronize, והתחביר <<<>>> - כתבנו מול השכבה הזו. מה שהופך אותו לנוח כל כך הוא בדיוק מה שהוא מסתיר מכם:

  • ניהול context מרומז. אינכם קוראים לשום פונקציית אתחול. בקריאת ה-Runtime הראשונה שנוגעת ב-GPU (למשל cudaMalloc או cudaSetDevice), ה-Runtime יוצר עבורכם באופן עצל (lazily) את ה-context של הdevice ומאתחל אותו. זהו ה-primary context, ונרחיב עליו בסעיף נפרד.
  • טעינת קוד אוטומטית. ה-fatbin (הבינארי השמן שראינו בשיעור 4.3) מוטמע בתוך קובץ ההרצה, וה-Runtime טוען ממנו את ה-SASS (אסמבלי החומרה) המתאים ל-GPU שלכם באופן שקוף, בלי שתקראו ל-cuModuleLoad בעצמכם.
  • התחביר <<<>>>. תחביר הlaunch המשולש הוא סוכר תחבירי (syntactic sugar) של CUDA C++ שהcompiler nvcc מתרגם לרצף קריאות Runtime (cudaConfigureCall / cudaLaunchKernel בגרסאות המודרניות). הוא זמין רק דרך ה-Runtime, לא דרך ה-Driver.

הנה תזכורת קצרה - חיבור וקטורים מלא מול ה-Runtime API, עם מאקרו בדיקת השגיאות CUDA_CHECK שמלווה אותנו מתחילת הקורס:

#include <cstdio>
#include <cstdlib>

#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)

__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);
    float *ha = (float*)malloc(bytes), *hb = (float*)malloc(bytes), *hc = (float*)malloc(bytes);
    for (int i = 0; i < n; ++i) { ha[i] = 1.0f; hb[i] = 2.0f; }

    float *da, *db, *dc;
    CUDA_CHECK(cudaMalloc(&da, bytes));            // <- the primary context is created here, quietly
    CUDA_CHECK(cudaMalloc(&db, bytes));
    CUDA_CHECK(cudaMalloc(&dc, bytes));
    CUDA_CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));

    int block = 256, grid = (n + block - 1) / block;
    vecAdd<<<grid, block>>>(da, db, dc, n);        // <- the Runtime's syntactic sugar
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));

    printf("c[0] = %.1f (expected 3.0)\n", hc[0]);
    cudaFree(da); cudaFree(db); cudaFree(dc);
    free(ha); free(hb); free(hc);
    return 0;
}

קישור - סטטי או דינמי. נקודה מעשית חשובה שתחזור בתרגול: את ה-Runtime API מותר לקשר סטטית אל תוך קובץ ההרצה (מוטמע ישירות ב-executable), וזו למעשה ברירת המחדל של nvcc (--cudart=static). לכן ldd על בינארי רגיל שקומפל ב-nvcc לא יראה את libcudart.so - הוא מוטמע פנימה. אפשר לבחור קישור דינמי במפורש עם --cudart shared, ואז libcudart.so יופיע. מסגרות (frameworks) כמו PyTorch נוהגות דווקא לטעון את libcudart.so דינמית. שני מודלי הטעינה חוקיים; ההיתר לקישור סטטי נובע מרישיון ה-EULA של ה-Toolkit (נספח A). וזכרו את הא-סימטריה: את ה-Runtime מותר לקשר סטטית, אבל את ה-Driver - כפי שנראה מיד - אף פעם לא.

ה-Driver API - הinterface הנמוך

ה-CUDA Driver API הוא הרובד ה-userspace של דרייבר NVIDIA - הרמה הנמוכה יותר שעליה ה-Runtime בנוי. הפונקציות שלו נושאות את הקידומת cu (למשל cuMemAlloc, cuInit, cuLaunchKernel), והוא ממומש בספרייה libcuda.so. בניגוד ל-Runtime, כאן אין קסמים מרומזים: אתם מנהלים בעצמכם את האתחול, את ה-context, ואת טעינת קוד ה-kernel. רצף העבודה הקנוני מול ה-Driver API הוא שרשרת מפורשת של שלבים:

פונקציה תפקיד
cuInit(0) אתחול ה-Driver API. חובה לקרוא פעם אחת לפני כל שאר ה-cu*.
cuDeviceGet(&dev, 0) השגת handle ל-GPU מספר 0.
cuCtxCreate(&ctx, 0, dev) יצירת context מפורש על הdevice ודחיפתו למחסנית ה-contexts של ה-thread.
cuModuleLoad(&mod, "k.ptx") טעינת module (PTX או cubin) מקובץ. cuModuleLoadData טוען מאותו image שכבר בזיכרון.
cuModuleGetFunction(&fn, mod, "vecAdd") חילוץ handle לפונקציית ה-kernel לפי שמה.
cuLaunchKernel(fn, gx,gy,gz, bx,by,bz, shmem, stream, args, extra) הlaunch של ה-kernel - המקבילה המפורשת של <<<>>>.

שימו לב במיוחד לשני שלבים שאין להם מקבילה גלויה ב-Runtime. cuModuleLoad / cuModuleLoadData הוא הטעינה הידנית של קוד ה-GPU: מה שה-Runtime עשה אוטומטית מה-fatbin המוטמע, כאן אתם עושים במפורש, ומתוך מקור שאתם בוחרים - קובץ .ptx, קובץ .cubin, מחרוזת PTX שנוצרה בזמן ריצה (בדיוק כפי שעשינו עם nvrtc בשיעור 4.3), או אפילו PTX שהורדתם מהרשת. cuLaunchKernel מקבל את הארגומנטים ל-kernel כמערך של מצביעים (void** args), ולא כרשימת פרמטרים מסודרת עם בדיקת טיפוסים בזמן compilation - זו אחת הנקודות שבהן ה-Driver API "בטיחותי" פחות אך גמיש יותר. שימו לב גם שהמאקרו <<<>>> פשוט לא זמין כאן; הוא סוכר של CUDA C++ מעל ה-Runtime בלבד.

קוד סגור. ה-Driver API הוא closed source. קיימים ניסיונות צד-שלישי לממש אותו מחדש (LibreCuda, וכן העבודה של tinygrad), אך המימוש הרשמי סגור.

קישור דינמי בלבד. בניגוד ל-Runtime, את ה-Driver API לא מקשרים סטטית - הוא תמיד מקושר דינמית אל libcuda.so. הסיבה עמוקה ומובילה אותנו ישר לסעיף התאימות קדימה: libcuda.so אינו חלק מה-Toolkit שאיתו קימפלתם, אלא חלק מהדרייבר המותקן על מכונת היעד. חייבים לקשר אליו דינמית כדי שהתוכנית תיקשר בזמן ריצה אל הדרייבר של אותה מכונה ספציפית, שיכול להיות חדש בהרבה מזה שקימפלתם מולו.

Context ראשי מול contexts מפורשים - primary vs explicit contexts

ה-context של CUDA הוא המקבילה על ה-GPU לתהליך (process) על ה-CPU: הוא מחזיק את מרחב הכתובות של הdevice, את ההקצאות, את ה-modules הטעונים ואת ה-streams. כאן טמון ההבדל ההתנהגותי הראשון והחשוב בין שתי הinterfaces, זה שה-glossary מפנה אליו במפורש כ"caveat סביב ניהול context".

  • Primary context - ה-context הראשי. לכל device יש context ראשי אחד, משותף, שה-Runtime API יוצר ומנהל עבורכם באופן מרומז. כשאתם קוראים ל-cudaMalloc הראשון, ה-Runtime מאחורי הקלעים "שולף" (retains) את ה-primary context של הdevice ומשתמש בו. כל קריאות ה-Runtime חולקות את ה-context היחיד הזה.
  • Explicit context - context מפורש. ה-Driver API מאפשר גם ליצור context משלכם עם cuCtxCreate. זהו context עצמאי, נפרד מה-primary, שאתם דוחפים ומוציאים ממחסנית ה-contexts של ה-thread בעצמכם (cuCtxPushCurrent / cuCtxPopCurrent). זה נותן בידוד (isolation) ושליטה עדינה - למשל כמה contexts שונים על אותו GPU.

מכאן נובע caveat אינטרופ (interop) חשוב: אם אתם מערבבים קריאות Driver API וקריאות Runtime API באותו thread, שני העולמות חייבים לחלוק את אותו context - אחרת הקצאה שעשיתם בצד אחד לא תהיה נראית בצד השני. הדרך הנכונה לערבב היא לא ליצור context מפורש עם cuCtxCreate, אלא לשלוף את ה-primary context עם cuDevicePrimaryCtxRetain - בדיוק אותו context שה-Runtime משתמש בו - וכך שתי הinterfaces מדברים על אותו מרחב זיכרון:

CUdevice dev;
cuDeviceGet(&dev, 0);
CUcontext primary;
cuDevicePrimaryCtxRetain(&primary, dev);   // the same context the Runtime will use
cuCtxSetCurrent(primary);                   // from now on cudaMalloc and cuMemAlloc share an address space
// ... safe to mix cu* and cuda* ...
cuDevicePrimaryCtxRelease(dev);

המסקנה המנטלית: ה-Runtime תמיד עובד על ה-primary context; ה-Driver יכול לעבוד עליו (עם PrimaryCtxRetain) או על context מפורש משלו (cuCtxCreate). זהו לב ההבדל בניהול ה-context בין שתי הinterfaces.

אותו kernel בשתי דרכים - same kernel, two ways

הדרך הטובה ביותר להרגיש את ההבדל היא להריץ את אותו kernel בדיוק פעם דרך ה-Runtime ופעם דרך ה-Driver. את גרסת ה-Runtime כבר ראינו למעלה. עכשיו נבנה את גרסת ה-Driver: נהדר את ה-kernel ל-PTX בנפרד, ואז נטען אותו ידנית ונשיק אותו. שימו לב שה-kernel חייב להיות extern "C" כדי שנוכל למצוא אותו לפי שם פשוט, בלי עיוות השמות (name mangling) של C++.

תחילה, קובץ ה-kernel לבדו, והcompilation שלו ל-PTX (ארכיטקטורה וירטואלית compute_90 של Hopper):

// kernel.cu  -  just the kernel, no main
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];
}

nvcc -arch=compute_90 -ptx kernel.cu -o kernel.ptx   # CUDA C++ -> PTX only

עכשיו תוכנית ה-host שטוענת את ה-PTX הזה דרך ה-Driver API בלבד. שימו לב שאין כאן ולו cuda* אחד - הכל cu*, ואנחנו קוראים במפורש את השלבים שה-Runtime היה עושה בשקט:

#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)

// read a text file (the PTX) into a string ending in '\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';   // important: 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);

    // ---- the explicit Driver API chain ----
    CU_CHECK(cuInit(0));
    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 to SASS
    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));  // the equivalent of <<<grid,block>>>
    CU_CHECK(cuCtxSynchronize());

    CU_CHECK(cuMemcpyDtoH(hc.data(), dc, bytes));
    printf("c[0] = %.1f (expected 3.0)\n", hc[0]);

    cuMemFree(da); cuMemFree(db); cuMemFree(dc);
    cuModuleUnload(mod); cuCtxDestroy(ctx);
    return 0;
}

g++ driver_vecadd.cpp -o driver_vecadd -I/usr/local/cuda/include -lcuda
./driver_vecadd

c[0] = 3.0 (expected 3.0)

שימו לב לשלוש נקודות. ראשית, שתי התוכניות מייצרות בדיוק אותה תוצאה ומריצות בדיוק אותו kernel על אותה חומרה - ההבדל הוא רק בשכבה שדרכה עברנו. שנית, גרסת ה-Driver ארוכה בהרבה: כל מה שהיה שורה אחת (cudaMalloc, <<<>>>) הפך לרצף מפורש. שלישית, שימו לב לדגל הקישור -lcuda (ולא -lcudart) - אנחנו נקשרים ישירות אל libcuda.so. הרגע שבו cuModuleLoadData רץ הוא בדיוק הרגע שבו ה-hardware driver מבצע JIT מ-PTX ל-SASS של sm_90, כפי שהסברנו בשיעור 4.3.

תאימות קדימה של ה-Driver API - forward compatibility

עכשיו נוכל להבין את אחת התכונות התפעוליות החשובות ביותר של ה-Driver API, זו שמסבירה מדוע הוא חייב להיות מקושר דינמית. ההבטחה מנוסחת כך: יישום שקומפל מול גרסה ישנה של ה-Driver API יכול לרוץ על מערכת עם גרסה חדשה יותר של ה-Driver API. במילים אחרות, אתם מקמפלים היום מול CUDA 12.0, שולחים בינארי ללקוח, והוא ירוץ בעוד שנתיים על מכונה עם דרייבר CUDA 13 מבלי שתקמפלו מחדש דבר.

המנגנון הוא שילוב של שלושה דברים שכבר פגשנו:

   build time (you)                            run time (at the client)

   the application links dynamically   ──sent──►   the loader loads libcuda.so
   against the old libcuda.so                       *installed* (newer) on the client
   (only records the dependency by name)             and links the binary to it
                                                     │
                                                     ▼
                                        the embedded PTX goes through JIT to SASS
                                        of the new GPU at load time
  1. קישור דינמי. הבינארי לא מכיל את libcuda.so בתוכו, אלא רק רשומת תלות בשם הזה. ה-loader של מערכת ההפעלה קושר את השם הזה, בזמן ריצה, אל ה-libcuda.so המותקן על מכונת היעד - שיכול להיות חדש בהרבה.
  2. תאימות ABI קדימה. NVIDIA מבטיחה שinterface ה-libcuda.so (ה-ABI - הinterface הבינארי) נשאר תואם קדימה: כל סמל (symbol) שהיה קיים בגרסה ישנה קיים גם בחדשה, עם אותה סמנטיקה. לכן בינארי ישן מוצא את כל מה שהוא צריך בספרייה החדשה.
  3. תאימות קדימה של PTX דרך JIT. אם הבינארי נושא PTX (ולא רק SASS), אזי גם ה-GPU עצמו יכול להיות מדור חדש יותר: ה-hardware driver יבצע JIT מה-PTX המוטמע אל ה-SASS של החומרה החדשה בזמן הטעינה, בדיוק כמו שראינו בשיעורים 0.4 ו-4.3.

הבחנה שחשוב לדייק בה: תאימות ה-libcuda.so היא קדימה - אפליקציה ישנה על דרייבר חדש. הכיוון ההפוך (אפליקציה שקומפלה מול Toolkit חדש, על דרייבר ישן) אינו מובטח באותו אופן, כי היא עלולה לדרוש סמל שעדיין לא היה קיים - ומכאן שגיאות מוכרות כמו CUDA driver version is insufficient for CUDA runtime version. השורה התחתונה לפריסה (deployment): מקמפלים מול ה-Toolkit הישן ביותר שאתם מתחייבים אליו, כוללים PTX כרשת ביטחון, ונשענים על כך שדרייברים חדשים יריצו את הבינארי הישן שלכם.

מתי באמת צריך את ה-Driver API - when you actually need it

אחרי כל זה, מתי כדאי לרדת ל-Driver API? התשובה הכנה: כמעט אף פעם, אם אתם כותבים יישום רגיל. מעט מאוד תוכניות נכתבות ישירות מולו, וה-Runtime API נותן לכם את אותה יכולת חישוב בפחות שורות ובלי מלכודות context. אבל יש חבורה מוגדרת של מקרים שבהם ה-Driver API הוא הכלי הנכון, ולעיתים היחיד:

  • טעינת קוד GPU שאינו ידוע בזמן build. אם ה-kernel נוצר בזמן ריצה (JIT של framework, מנוע inference שמייצר kernels לפי הקלט, קוד שהתקבל מהרשת) - אתם חייבים cuModuleLoadData כדי לטעון PTX/cubin שאינו מוטמע ב-fatbin. זה בדיוק מה שעשינו עם nvrtc בשיעור 4.3, ולכן כלל ה-frameworks שמייצרים קוד דינמית (CuPy, מנועי DL מסוימים) נשענים על ה-Driver API.
  • ניהול context מפורש ומרובה. כשצריכים כמה contexts מבודדים על אותו GPU, שליטה עדינה על מחסנית ה-contexts, או ניהול ריבוי-GPU ידני מדויק - הבקרה המפורשת של ה-Driver API הכרחית.
  • הפחתת תלות בזמן פריסה. תוכנית שנקשרת רק ל-libcuda.so (שמגיע עם הדרייבר, קיים תמיד על מכונה עם GPU) ולא ל-libcudart.so נהנית מתלות מינימלית - שיקול פריסה עבור כלים שרוצים "רק דרייבר, בלי Toolkit".
  • כלים ומסגרות שבונים מעל CUDA. מי שמממש runtime משלו, פרופיילר, או שכבת שפה חדשה מעל ה-GPU - עובד ברמת ה-Driver API כי הוא צריך את הבקרה שה-Runtime מסתיר.

הכלל האצבע: אם אתם שואלים את עצמכם "האם אני צריך את ה-Driver API?", כנראה שלא. אתם תדעו שאתם צריכים אותו כשתיתקלו בקיר קונקרטי - kernel שנוצר בזמן ריצה, context שצריך בידוד, או תלות שצריך לגזום - שה-Runtime API פשוט לא נותן לכם לטפס עליו.

משמעת שמות ואבחון מהיר - naming discipline and diagnostics

נסגור עם משמעת השמות שתשמור אתכם מחצי מבלבולי הפרק, ועם אבחון מעשי שתשתמשו בו בתרגול. ראשית, המילון:

  • cu* = פונקציית Driver API. cuda* = פונקציית Runtime API.
  • libcuda.so = מימוש ה-Driver API, מגיע עם הדרייבר.
  • libcudart.so = מימוש ה-Runtime API, מגיע עם ה-Toolkit (שימו לב ל-rt).
  • nvidia.ko = מודול הליבה (kernel mode), מתחת לשתיהן - נעסוק בו בשיעור 5.2.

שנית, האבחון. הכלי ldd מראה לאילו ספריות דינמיות בינארי נקשר, וכך אפשר לזהות בעין את השכבות. שימו לב לשלושה תרחישים:

nvcc vecadd.cu -o vecadd                 # default: cudart *static*
ldd vecadd | grep -E 'cuda|cudart'

        libcuda.so.1 => /lib/x86_64-linux-gnu/libcuda.so.1

בברירת המחדל ה-Runtime מקושר סטטית ולכן libcudart אינו מופיע - הוא מוטמע פנימה. אבל libcuda.so.1 (הדרייבר) כן מופיע: אפילו תוכנית Runtime טהורה נקשרת בסופו של דבר לדרייבר, כי ה-Runtime עוטף אותו. עכשיו נכריח קישור דינמי של ה-Runtime:

nvcc --cudart shared vecadd.cu -o vecadd_shared
ldd vecadd_shared | grep -E 'cuda|cudart'

        libcudart.so.12 => /usr/local/cuda/lib64/libcudart.so.12
        libcuda.so.1 => /lib/x86_64-linux-gnu/libcuda.so.1

עכשיו רואים את שתי השכבות בבירור: libcudart.so.12 (Runtime, מה-Toolkit) מעל libcuda.so.1 (Driver, מהדרייבר). ולבסוף, תוכנית ה-Driver הטהורה שלנו נקשרת רק ל-libcuda.so.1, בלי cudart כלל.

אבחון תפעולי חשוב: אם libcuda.so חסר מהמערכת, המשמעות כמעט תמיד היא שהדרייבר של ה-GPU לא הותקן כראוי (בניגוד ל-libcudart.so החסר, שמשמעו רק שה-Toolkit לא מותקן או לא ב-path). זו בדיקת השפיות הראשונה כשתוכנית CUDA לא עולה: יש libcuda.so.1? אם לא, התקינו דרייבר.

סיכום

  • שתי הinterfaces שהיישום מדבר דרכם ל-GPU הם ה-CUDA Runtime API (גבוה, cuda*, libcudart.so, הנפוץ) וה-CUDA Driver API (נמוך, cu*, libcuda.so); ה-Runtime עוטף את ה-Driver ויושב מעליו.
  • ה-Runtime API מנהל context באופן מרומז, טוען את קוד ה-kernel אוטומטית מה-fatbin, ומספק את התחביר <<<>>>; מותר לקשר אותו סטטית (ברירת המחדל של nvcc) או דינמית.
  • ה-Driver API מפורש: cuInit, cuDeviceGet, cuCtxCreate, cuModuleLoad/cuModuleLoadData, cuModuleGetFunction, cuLaunchKernel; הוא קוד סגור, ומקושר דינמית בלבד אל libcuda.so.
  • ה-Runtime תמיד עובד על ה-primary context המשותף של הdevice; ה-Driver יכול לעבוד עליו (cuDevicePrimaryCtxRetain, נדרש כדי לערבב עם Runtime) או ליצור context מפורש ומבודד (cuCtxCreate).
  • אותו kernel רץ זהה בשתי הדרכים; גרסת ה-Driver ארוכה בהרבה כי כל מה שהיה מרומז ב-Runtime הופך לשלב מפורש, והיא נקשרת עם -lcuda.
  • תאימות קדימה: בינארי שקומפל מול Driver API ישן ירוץ על דרייבר חדש יותר, בזכות קישור דינמי ל-libcuda.so, תאימות ABI קדימה, ו-JIT של PTX ל-SASS על חומרה חדשה; הכיוון ההפוך (Toolkit חדש על דרייבר ישן) אינו מובטח.
  • כמעט תמיד כותבים מול ה-Runtime; יורדים ל-Driver API רק כשצריך לטעון kernel שנוצר בזמן ריצה, לנהל contexts מבודדים, לגזום תלות פריסה, או לבנות framework מעל CUDA.
  • משמעת שמות: cu*=Driver, cuda*=Runtime, libcuda.so=Driver (מהדרייבר), libcudart.so=Runtime (מה-Toolkit); libcuda.so חסר משמעו שהדרייבר לא הותקן.
  • אבחון עם ldd: ברירת המחדל של nvcc מסתירה את libcudart (סטטי) אך תמיד מראה libcuda.so.1; --cudart shared חושף את שתי השכבות; את פרטי הדרייבר עצמו שמתחת לשתיהן נפרק בשיעור 5.2.