לדלג לתוכן

5.1 Driver API מול Runtime API תרגול

תרגול - Driver API מול Runtime API

בתרגול הזה נרגיש בידיים את שתי הinterfaces שהכרנו בהרצאה. נתחיל מ-baseline מוכר של ה-Runtime API, נעבור לפרק את אותו kernel בדיוק דרך ה-Driver API הנמוך, נשתמש ב-ldd כדי לזהות בעין את שכבות המחסנית בתוך בינארי מהודר, ולבסוף ננתח את הבטחת התאימות קדימה ונחליט מתי באמת צריך לרדת ל-Driver API. עבדו לפי הסדר - כל תרגיל נשען על הקודם. כל התרגילים רצים על H100 (או כל GPU של NVIDIA, כולל T4 בחינם ב-Google Colab); היכן שרלוונטי, החליפו את sm_90 / compute_90 בארכיטקטורה של הכרטיס שלכם.

הכנה

צרו תיקיית עבודה, ובתוכה שני קבצים. הראשון הוא ה-kernel לבדו, מסומן extern "C" כדי שנוכל למצוא אותו לפי שם דרך ה-Driver API:

// kernel.cu
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];
}

השני הוא תוכנית Runtime API מלאה שמשתמשת באותו kernel (העתיקו את ה-kernel גם לתוכה, או #include "kernel.cu"):

// runtime_vecadd.cu
#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)

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);
    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));
    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);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));

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

תרגיל 1 - baseline של ה-Runtime API

הריצו את התוכנית המוכרת כדי לקבע נקודת ייחוס לפני שנרד לרובד הנמוך.

  1. הדרו את runtime_vecadd.cu עם nvcc -arch=sm_90 runtime_vecadd.cu -o runtime_vecadd.
  2. הריצו ./runtime_vecadd וודאו שהתוצאה היא 3.0.
  3. סמנו לעצמכם: כמה קריאות cuda* היו דרושות מרגע ההקצאה ועד הlaunch? היכן, בדיוק, נוצר ה-context של הdevice?

רמז: אין בקוד שום קריאת אתחול מפורשת. ה-primary context נוצר בשקט בקריאה הראשונה שנוגעת ב-GPU. איזו קריאה זו?

תרגיל 2 - זיהוי השכבות עם ldd

עכשיו נזהה בעין את libcudart מול libcuda בתוך הבינארי, ונבין את השפעת מודל הקישור.

  1. הריצו ldd runtime_vecadd | grep -E 'cuda|cudart'. האם libcudart מופיע? האם libcuda מופיע? הסבירו את מה שאתם רואים לאור ברירת המחדל של nvcc.
  2. הדרו מחדש עם קישור Runtime דינמי מפורש: nvcc --cudart shared -arch=sm_90 runtime_vecadd.cu -o runtime_shared, והריצו שוב ldd runtime_shared | grep -E 'cuda|cudart'. מה השתנה?
  3. הסבירו במשפט אחד לכל שורה: מאיזו חבילת התקנה מגיע libcudart.so, ומאיזו מגיע libcuda.so.

רמז: ברירת המחדל של nvcc היא --cudart=static. ספרייה שמקושרת סטטית מוטמעת בתוך הבינארי ולכן לא תופיע ב-ldd. שימו לב שאפילו כשה-Runtime סטטי, שכבה אחת עדיין מופיעה תמיד - מדוע?

תרגיל 3 - אותו kernel דרך ה-Driver API

זה לב התרגול: נריץ את אותו vecAdd בדיוק, אבל דרך ה-Driver API בלבד, מתוך PTX שנטען ידנית.

  1. הדרו את ה-kernel ל-PTX בלבד: nvcc -arch=compute_90 -ptx kernel.cu -o kernel.ptx. הציצו בקובץ - זהו טקסט PTX קריא.
  2. כתבו driver_vecadd.cpp שמבצע את השרשרת המפורשת: cuInit -> cuDeviceGet -> cuCtxCreate -> קריאת kernel.ptx לזיכרון -> cuModuleLoadData -> cuModuleGetFunction("vecAdd") -> cuMemAlloc + cuMemcpyHtoD -> cuLaunchKernel -> cuCtxSynchronize -> cuMemcpyDtoH. השתמשו במאקרו CU_CHECK כמו בהרצאה ובשיעור 4.3.
  3. הדרו וקשרו מול הדרייבר: g++ driver_vecadd.cpp -o driver_vecadd -I/usr/local/cuda/include -lcuda. הריצו וודאו שהתוצאה זהה ל-3.0.
  4. הריצו ldd driver_vecadd | grep -E 'cuda|cudart'. במה רשימת הספריות שונה מזו של runtime_vecadd?

רמז: את הארגומנטים ל-cuLaunchKernel מעבירים כמערך void* args[] = { &da, &db, &dc, &n }. שימו לב שקוראים את קובץ ה-PTX עם תו סיום \0 - cuModuleLoadData מצפה ל-image מסתיים ב-null. הפונקציה נמצאת לפי השם "vecAdd" בזכות ה-extern "C".

תרגיל 4 - הבטחת התאימות קדימה

כאן לא נכתוב קוד אלא ננתח, ונאסוף ראיות מהבינארים שכבר בנינו.

  1. נסחו במילים שלכם את הבטחת התאימות קדימה של ה-Driver API: אילו שני צדדים (אפליקציה, דרייבר) ישנים ואילו חדשים בתרחיש המובטח, ומה לא מובטח בכיוון ההפוך.
  2. הסבירו את שלושת המנגנונים שמאפשרים את ההבטחה: קישור דינמי ל-libcuda.so, תאימות ABI קדימה, ו-JIT של PTX. לכל מנגנון, הצביעו על הראיה בבינארי או בפקודות שכבר הרצתם (למשל: איזו פקודת ldd מוכיחה את הקישור הדינמי? איזו פקודת cuobjdump על הבינארי מוכיחה שיש PTX שיכול לעבור JIT קדימה?).
  3. הריצו nvidia-smi והסתכלו על שדה ה-CUDA Version בפינה הימנית העליונה מול פלט של nvcc --version. הסבירו איזה מהם הוא "גרסת הדרייבר" ואיזה "גרסת ה-Toolkit", ואיך זה קשור לתאימות קדימה.

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

תרגיל 5 - מתי צריך את ה-Driver API

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

  1. מנוע inference שמקבל מהמשתמש נוסחה מתמטית, מייצר ממנה קוד CUDA C++ כמחרוזת בזמן ריצה, ורוצה להריץ אותה על ה-GPU.
  2. תוכנית חישוב מדעי סטנדרטית שמכפילה שתי מטריצות בגודל קבוע הידוע בזמן compilation.
  3. כלי שרוצה להיפרש (deploy) על מכונות ענן עם דרייבר בלבד, בלי להתקין את ה-CUDA Toolkit כלל, ולכן חייב תלות מינימלית.
  4. framework שרוצה ליצור שני contexts מבודדים לחלוטין על אותו GPU כדי לבודד שני משתמשים.

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

תרגיל 6 (בונוס) - אינטרופ בין שתי הinterfaces

נאמת בעצמנו את ה-caveat של ניהול ה-context: מה קורה כשמערבבים cuda* ו-cu*.

  1. כתבו תוכנית קטנה שקוראת cuInit(0), cuDeviceGet, ואז cuDevicePrimaryCtxRetain + cuCtxSetCurrent כדי לשלוף את ה-primary context של הdevice.
  2. באותה תוכנית, הקצו זיכרון עם cudaMalloc (Runtime) והשיקו את vecAdd דרך cuLaunchKernel (Driver) על אותו מצביע. ודאו שזה עובד - שתי הinterfaces חולקים את ה-context.
  3. עכשיו החליפו את cuDevicePrimaryCtxRetain ב-cuCtxCreate (context מפורש חדש, לא ה-primary), והריצו שוב. הסבירו מה משתבש עקרונית כשה-Runtime וה-Driver מפוצלים לשני contexts שונים (רמז: הקצאות אינן נראות בין contexts).

רמז: ה-Runtime תמיד עובד על ה-primary context. כדי לערבב אותו עם ה-Driver בבטחה, ה-Driver חייב לעבוד על אותו primary context (cuDevicePrimaryCtxRetain), ולא על context עצמאי שיצרתם ב-cuCtxCreate.