לדלג לתוכן

5.3 Runtime API, streams ו events הרצאה

בשיעור 0.4 מיפינו את שרשרת התוכנה שבין הקוד שלנו לבין החומרה, ובשיעור 3.2 פגשנו לראשונה את המושגים pinned, cudaMemcpyAsync, streams ו-events - אבל הצגנו אותם שם ברמת המבוא ההכרחית בלבד כדי לחפוף העברות, ודחינו במפורש את הטיפול המלא לשיעור הזה. עכשיו אנחנו חוזרים אל ה-CUDA Runtime API (interface זמן-הריצה של CUDA) ומעמיקים בשלושת העמודים שעליהם נשענת כל תוכנית CUDA רצינית: מודל השגיאות, שבלעדיו אתם עובדים באפלה; ה-streams (זרמים) לעומק, כולל ההבחנה העדינה והחשובה בין ה-default stream המסורתי (legacy) לבין per-thread default stream; וה-events (אירועים) ככלי למדידת זמן GPU מדויקת ולבניית תלויות בין זרמים. המסר המרכזי: מעל ה-Runtime API בונים לא רק קוד שרץ, אלא קוד שרץ במקביל - חופף העתקות עם חישוב - ונכשל ברעש במקום להחזיר תשובה שגויה בשקט. את הכלים האלה נשתמש שוב בשיעור 5.5 (CUDA Graphs) ובפרק 9 (profiling), אז שווה לבסס אותם היטב כאן.

מודל השגיאות של ה-Runtime - the error model

כמעט כל פונקציה ב-CUDA Runtime API מחזירה ערך מסוג cudaError_t - מספר שלם שבו cudaSuccess (הערך 0) פירושו הצלחה וכל ערך אחר הוא קוד שגיאה. את המחרוזת הקריאה מפיקים עם cudaGetErrorString(err), ואת השם הסמלי (למשל cudaErrorInvalidValue) עם cudaGetErrorName(err). הבעיה: הרבה מהפונקציות הן אסינכרוניות מנקודת מבט הhost (host, כלומר ה-CPU), ובראשן launch ה-kernel, שאין לה בכלל ערך החזרה כי ה-<<< >>> איננו קריאת פונקציה רגילה. לכן ה-Runtime מנהל, בנוסף לערכי ההחזרה, גם שגיאה אחרונה per-thread שאליה אפשר לגשת בשתי דרכים:

פונקציה מחזירה מאפסת?
cudaPeekAtLastError() את השגיאה האחרונה שנרשמה ב-thread הזה לא - רק מציצה
cudaGetLastError() את אותה שגיאה כן - מאפסת חזרה ל-cudaSuccess

ההבחנה מהותית. אם אתם רק רוצים לבדוק "האם קרתה שגיאה" בלי לצרוך אותה, השתמשו ב-cudaPeekAtLastError. אם אתם רוצים לבדוק וגם לנקות את הדגל כדי שהבדיקה הבאה תתחיל מדף חלק, השתמשו ב-cudaGetLastError. תבנית נפוצה מאוד: קריאה יחידה ל-cudaGetLastError() בתחילת קטע קריטי כדי "לבלוע" שגיאה ישנה שאולי נשארה תלויה, לפני שמתחילים למדוד קטע חדש.

שגיאות סינכרוניות מול אסינכרוניות בlaunch kernel. כשאתם משיקים kernel יש שני סוגי שגיאות שונים לחלוטין, וזו נקודת המכשול הנפוצה ביותר:

  • שגיאת launch (synchronous / launch error). נתגלית מיד בעת הlaunch: תצורה בלתי-חוקית, למשל יותר threads לבלוק ממה שהחומרה מרשה (מעל 1024), יותר shared memory דינמי ממה שיש, או grid בגודל 0. שגיאה כזו נרשמת בדגל ה"שגיאה האחרונה" ותיקָּרֵא ב-cudaGetLastError() מיד אחרי ה-<<< >>>.
  • שגיאת הרצה (asynchronous / execution error). קורית בזמן שה-kernel רץ על ה-GPU, זמן רב אחרי שהhost כבר המשיך הלאה: גישה מחוץ לתחום ב-global memory, assert שנכשל, חלוקה שגרמה ל-trap. שגיאה כזו לא יכולה להיתפס מיד - היא צצה רק בקריאה המסנכרנת הבאה: cudaDeviceSynchronize, cudaStreamSynchronize, או אפילו cudaMemcpy שממתין לתוצאה.

מכאן נובע כלל הזהב: כדי לבדוק kernel במלואו צריך שתי בדיקות. הראשונה מיד אחרי הlaunch תופסת שגיאות תצורה; השנייה, אחרי סנכרון, תופסת שגיאות הרצה.

myKernel<<<grid, block>>>(d_a, d_b, n);
CUDA_CHECK(cudaGetLastError());        // catches a launch error (config)
CUDA_CHECK(cudaDeviceSynchronize());   // catches an execution error (out-of-bounds etc.)

בקוד production משאירים לרוב את ה-cudaDeviceSynchronize הזה רק ב-build של debug (הוא עולה ביצועים), אך את cudaGetLastError() אחרי כל launch כדאי להשאיר תמיד - הוא כמעט חינם.

שגיאות דביקות ואתחול עצל - sticky errors and lazy init

לא כל השגיאות שוות. יש להבחין בין שני מעמדות:

  • שגיאה לא-דביקה (non-sticky). נובעת מבקשה לא-חוקית שה-Runtime דחה בלי לגעת ב-GPU: למשל cudaMalloc שנכשל ב-cudaErrorMemoryAllocation כי אין מספיק זיכרון, או cudaErrorInvalidValue על ארגומנט שגוי. ה-context (הֶקְשֵׁר) נשאר תקין; אתם קוראים ל-cudaGetLastError כדי לנקות את הדגל וממשיכים כרגיל.
  • שגיאה דביקה (sticky). נובעת מכך ש-kernel כבר רץ ועשה משהו קטלני - גישה לזיכרון לא-חוקי, misaligned address, trap של החומרה. שגיאה כזו משחיתה את ה-context כולו. מרגע זה, כל קריאת Runtime עוקבת - אפילו cudaMalloc תמים לגמרי לא קשור - תחזיר את אותה שגיאה דביקה, שוב ושוב. cudaGetLastError לא ינקה אותה.

הדרך היחידה להתאושש משגיאה דביקה היא להרוג את ה-context: או לסיים את התהליך, או לקרוא cudaDeviceReset() (שמשמיד את כל ההקצאות וה-streams וה-events ומאתחל מאפס). אין "לתפוס ולהמשיך" כמו ב-exception של C++. לכן שגיאה דביקה בקוד CUDA היא באג חמור שיש לתקן, לא מצב שמנהלים בזמן ריצה. זו בדיוק הסיבה שבודקים כל קריאה: ברגע שתופסים שגיאה דביקה, ההודעה הראשונה שתראו (עם קובץ ושורה) היא זו שמצביעה על מקור הבעיה; אם לא תבדקו, תראו את השגיאה "מתגלגלת" לקריאה אקראית מאוחרת ותבזבזו שעות.

אתחול עצל - lazy initialization. ה-Runtime איננו יוצר את ה-context בטעינת התוכנית, אלא בפעם הראשונה שאתם קוראים לפונקציית Runtime שדורשת אותו (הקצאה, launch וכדומה). לכן העלות של האתחול - זיהוי הdevice, טעינת מודולים, בניית ה-context - "נדבקת" לקריאה הראשונה שלכם ומזהמת מדידת זמן תמימה. תרגיל נפוץ הוא לאלץ אתחול מוקדם עם קריאה זולה כמו cudaFree(0) או cudaSetDevice(0) לפני שמתחילים למדוד. שכבה נוספת של עצלות היא טעינת מודולים עצלה (lazy module loading), שמאז CUDA 11.7 היא ברירת המחדל ונשלטת במשתנה הסביבה CUDA_MODULE_LOADING=LAZY: קוד ה-kernel (ה-SASS) נטען ל-GPU רק כשה-kernel מושק לראשונה, לא בעליית התוכנית. תוצאה מעשית: הlaunch הראשונה של כל kernel יקרה מעט יותר. לכן, כפי שהדגשנו בשיעור 3.2, תמיד מריצים איטרציית חימום (warm-up) לפני מדידה.

מאקרו לבדיקת שגיאות - a robust CUDA_CHECK macro

לכתוב if (err != cudaSuccess) { ... } אחרי כל קריאה זה בלתי-נסבל, ולכן כל בסיס קוד CUDA רציני מגדיר מאקרו יחיד שעוטף כל קריאה, בודק את קוד ההחזרה, ואם הוא שגוי - מדפיס קובץ, שורה, הקריאה המדויקת וההודעה, ואז יוצא. זהו המאקרו שנשתמש בו לכל אורך הקורס:

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CUDA_CHECK(call)                                                       \
    do {                                                                       \
        cudaError_t err_ = (call);                                            \
        if (err_ != cudaSuccess) {                                            \
            fprintf(stderr, "CUDA error at %s:%d\n  call: '%s'\n  -> %s (%s)\n",\
                    __FILE__, __LINE__, #call,                               \
                    cudaGetErrorString(err_), cudaGetErrorName(err_));       \
            exit(EXIT_FAILURE);                                              \
        }                                                                    \
    } while (0)

שימו לב לפרטים המרכיבים מאקרו נכון. ה-do { ... } while (0) הופך את המאקרו למשפט יחיד שמתנהג נכון גם בתוך if בלי סוגריים מסולסלים. המשתנה err_ עם קו תחתי מקטין סיכוי להתנגשות שם. ה-#call (סטרינגיפיקציה) מדפיס את הטקסט המקורי של הקריאה, כך שההודעה אומרת לכם בדיוק איזו שורה נכשלה. ולסיכום נוסיף מאקרו-אחות במיוחד לlaunches kernel, שמאגד את שתי הבדיקות מהסעיף הקודם:

// full check of a kernel launch (config + execution). Sync only in debug builds.
#define CUDA_CHECK_KERNEL()                                                    \
    do {                                                                       \
        CUDA_CHECK(cudaGetLastError());        /* launch error (config) */     \
        CUDA_CHECK(cudaDeviceSynchronize());   /* execution error */          \
    } while (0)

מדוע לבדוק כל קריאה, גם cudaEventDestroy תמים? כי שגיאות דביקות מתגלגלות. אם cudaFree מחזיר שגיאה, מקור הבעיה כמעט תמיד קדם לו - kernel שרץ קודם והשחית את ה-context. הודעה עם קובץ ושורה הופכת ציד באג של שעות לתיקון של דקות. העלות של הבדיקה זניחה (השוואה למספר), התועלת עצומה - זו אחת ההרגלים שמפרידים בין קוד CUDA חובבני למקצועי. ההערה היחידה: אין לקרוא ל-exit מתוך destructor או קוד שחייב לשחרר משאבים; בספריות מעדיפים לזרוק exception במקום exit, אך לקוד שלנו ולתרגילים - exit פשוט וברור.

זרמים לעומק - streams in depth

stream (זרם) הוא תור של פעולות GPU שמתבצעות לפי הסדר שבו הונפקו. חזרנו על שתי עובדות היסוד בשיעור 3.2, ונחדד אותן:

  • בתוך אותו stream הפעולות סדרתיות: פעולה i+1 לא מתחילה עד שפעולה i הסתיימה. זו ערובת התלות שעליה אתם סומכים.
  • בין streams שונים אין שום סדר מובטח; ה-GPU רשאי, ואף אמור, להריץ אותם במקביל, בכפוף למשאבי החומרה (SMs פנויים, Copy Engines פנויים).

מכאן שה-stream הוא המנוף לקבילות (concurrency): פיזור עבודה בלתי-תלויה על כמה streams מאפשר ל-GPU לחפוף חישוב עם חישוב, וחישוב עם העתקה. אבל - ואת זה 3.2 רק רמז - יש עדינות עמוקה בשאלה מהו ה-default stream, וזו נקודת הכשל הנפוצה בקוד קבילות.

ה-default stream המסורתי - legacy default stream (stream 0). בכל הקוד שכתבנו עד עכשיו, שלא ציין stream, הפעולות רצו ב-default stream. ב-semantics המסורתי, ה-default stream הוא מסנכרן (implicitly synchronizing): פעולה בו חוסמת את כל ה-streams ה"רגילים" (blocking streams) על הdevice, וכל אחד מהם חוסם אותו. התוצאה ההרסנית: אם תשתלשל ולו פעולה אחת אל ה-default stream באמצע צנרת של streams מפורשים, היא תיצור מחסום שמבטל את כל הקבילות שבניתם.

   legacy default stream (stream 0) synchronizes everyone:

   stream A:  ──K_A──┐
   default :        [ M ]           ← operation on default
   stream B:  ──K_B──┘   └──K_B'──   ← B waits for default, and default waits for A and B
              ▲ K_A and K_B cannot overlap M - everything gets serialized

ה-per-thread default stream. מאז CUDA 7 יש התנהגות חלופית ל-default stream, שנבחרת בזמן הcompilation בדגל nvcc --default-stream per-thread (או בהגדרת המאקרו CUDA_API_PER_THREAD_DEFAULT_STREAM לפני ה-#include). במצב הזה לכל host thread יש default stream משלו, וה-default stream הזה אינו מסנכרן מול streams אחרים - הוא מתנהג כמו stream רגיל ובלתי-חוסם. זהו שינוי משמעותי לתוכניות מרובות-threads: שני threads של הhost, כל אחד "סתם" משיק kernel בלי לציין stream, ירוצו כעת במקביל במקום להסתדרן זה מול זה.

היבט legacy default stream per-thread default stream
כמה default streams אחד לכל התהליך אחד לכל host thread
מסנכרן מול streams רגילים? כן (חוסם הכל) לא (בלתי-חוסם)
נבחר איך ברירת המחדל nvcc --default-stream per-thread
מתי חשוב תמיד קריטי בקוד מרובה-threads

streams חוסמים מול בלתי-חוסמים. גם כשיוצרים stream מפורש יש בחירה. cudaStreamCreate(&s) יוצר blocking stream - כזה שכן מסתנכרן מול ה-legacy default stream. cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) יוצר stream שאינו מסתנכרן מול ה-default stream כלל - מבודד לחלוטין. לצנרת קבילות אמיתית, שבה אתם רוצים שהזרמים יהיו עצמאיים ככל האפשר, ה-cudaStreamNonBlocking הוא הבחירה הבטוחה. יצירה, עדיפות והשמדה:

cudaStream_t s;
CUDA_CHECK(cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking));
// ... usage ...
CUDA_CHECK(cudaStreamDestroy(s));

// priorities: "important" streams take precedence in scheduling over "regular" streams
int lo, hi;
CUDA_CHECK(cudaDeviceGetStreamPriorityRange(&lo, &hi));  // hi < lo, hi = higher priority
cudaStream_t urgent;
CUDA_CHECK(cudaStreamCreateWithPriority(&urgent, cudaStreamNonBlocking, hi));

שימו לב שבטווח העדיפויות מספר קטן יותר = עדיפות גבוהה יותר (hi הוא הערך המספרי הקטן). העדיפות היא רמז לתזמון בין blocks שממתינים ל-SM, לא ערובה קשיחה.

פעולות אסינכרוניות - async operations on a stream

הכלי שהופך stream ממסגרת תיאורטית לצנרת אמיתית הוא ההעתקה האסינכרונית. cudaMemcpyAsync(dst, src, bytes, kind, stream) זהה ל-cudaMemcpy פרט לשני הבדלים: יש לה פרמטר stream, והיא חוזרת לhost מיד, בלי לחכות לסיום ההעברה. הlaunch מקבלת אף היא stream בשדה הרביעי של תצורת הlaunch, <<<grid, block, sharedBytes, stream>>> (השדה השלישי הוא shared memory דינמי, לרוב 0):

CUDA_CHECK(cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, s));
myKernel<<<grid, block, 0, s>>>(d_a, d_c, n);     // runs after the copy - same stream
CUDA_CHECK(cudaGetLastError());                   // checking a launch error
CUDA_CHECK(cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, s));
CUDA_CHECK(cudaStreamSynchronize(s));             // here the host really waits for everything to finish

שלושת הצעדים ב-stream s יתבצעו לפי הסדר (copy, kernel, copy) - זו ערובת התוך-stream - אך הhost לא נחסם עד ה-cudaStreamSynchronize בסוף. התנאי ההכרחי, שקל לשכוח ולכן נחזור עליו: כדי ש-cudaMemcpyAsync תהיה באמת אסינכרונית, זיכרון הhost (h_a, h_c) חייב להיות pinned (נעול-עמודים), כפי שלמדנו בשיעור 3.2. אם תעבירו מצביע pageable, הדרייבר נאלץ לעבור דרך bounce buffer ולסנכרן - ה"async" מתנוון בשקט ל-copy חוסם, בלי שום הודעת שגיאה. זהו הבאג הקלאסי "הקוד שלי משתמש ב-streams אבל לא מקבל שום חפיפה": כמעט תמיד הזיכרון pageable.

תזכורת מדוע החומרה מסוגלת לחפוף: ב-H100 מנועי ה-DMA (Copy Engines) עצמאיים לחלוטין מיחידות החישוב (SMs), ויש כמה מהם, כך שהכרטיס יכול לבצע העתקת H2D, חישוב, והעתקת D2H בו-זמנית - שלושה זרמי עבודה במקביל - אם רק תספקו לו את העבודה מפוזרת על streams שאינם ה-default.

מדידת זמן GPU עם events - precise timing

event (אירוע) הוא סמן שמניחים בתוך stream, ושמוקלט (נחתם בחותמת-זמן) ברגע שכל הפעולות שקדמו לו באותו stream הסתיימו. התפקיד הראשון של event הוא מדידת זמן GPU מדויקת. בניגוד ל-std::chrono שמודד על ה-CPU ותופס גם את overhead הסנכרון והלכלוך של הhost, event מוקלט על ה-GPU עצמו ומודד את הזמן שבו ה-GPU באמת עבד, ברזולוציה של כחצי מיקרו-שנייה:

cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));

// warm-up: one launch we don't measure, to absorb lazy init and JIT
myKernel<<<grid, block>>>(d_a, d_c, n);
CUDA_CHECK(cudaDeviceSynchronize());

CUDA_CHECK(cudaEventRecord(start));                  // mark before
myKernel<<<grid, block>>>(d_a, d_c, n);
CUDA_CHECK(cudaEventRecord(stop));                   // mark after
CUDA_CHECK(cudaEventSynchronize(stop));              // wait for stop to be recorded on the GPU
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));  // difference in milliseconds (float)
printf("kernel: %.3f ms\n", ms);

CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));

שלוש נקודות עדינות. ראשית, cudaEventRecord(start) בלי ציון stream מקליט ב-default stream; אם אתם מודדים עבודה ב-stream ספציפי, חובה להעביר אותו: cudaEventRecord(start, s), אחרת החותמת תילקח ביחס לזרם הלא-נכון. שנית, cudaEventElapsedTime דורש ששני ה-events כבר הוקלטו בפועל - ומכאן ה-cudaEventSynchronize(stop) שממתין לחתימת ה-stop לפני קריאת ההפרש; בלעדיו תקבלו cudaErrorNotReady. שלישית, המדידה מתקבלת ב-float מילישניות, אז לקטעים קצרצרים מודדים לולאה של הרבה חזרות ומחלקים, כדי לצאת מרעש הרזולוציה.

דוגמה מספרית מלאה. נניח שמדדנו kernel של חיבור וקטורים על n = 2^26 = 67,108,864 איברי float, וקיבלנו ms = 1.20. הגישות לזיכרון: קוראים שני מערכים וכותבים אחד, שלושה float (12 בתים) לכל איבר, סך הכל 67,108,864 x 12 = 805,306,368 בתים ≈ 0.805 GB. רוחב הפס האפקטיבי: 0.805 GB / (1.20 / 1000) s ≈ 671 GB/s. מול רוחב הפס של ה-HBM3 ב-H100, בסדר גודל של 3.35 TB/s, זה כ-20% - סימן שה-kernel memory-bound (memory-bound) ורחוק מהרוויה, בדיוק כפי שמצפים מ-vecAdd, שה-arithmetic intensity שלו זעומה. ה-event הוא הכלי שנתן לנו את המספר 1.20 הזה במדויק.

תלויות בין זרמים וסנכרון - cross-stream dependencies

התפקיד השני של event הוא בניית תלות בין streams בלי לחסום את הhost. cudaStreamWaitEvent(streamB, event, 0) מורה ל-streamB: אל תתחיל את הפעולות הבאות שלך עד שה-event (שאולי הוקלט ב-streamA) יתרחש. כך יוצרים תלות "צולבת" בין זרמים - למשל, kernel ב-streamB שצריך את התוצאה של kernel ב-streamA - בלי לשרשר את שניהם לאותו stream ובלי להעיר את ה-CPU.

   cudaStreamWaitEvent builds a cross-stream dependency:

   stream A:  ──K_A──►[ record e ]
                              │  (e is recorded when K_A finishes)
                              ▼
   stream B:  ─────────[ wait e ]──►──K_B──   ← K_B waits for e, the rest of B runs freely

cudaEvent_t e;
CUDA_CHECK(cudaEventCreateWithFlags(&e, cudaEventDisableTiming));  // for sync only
kernelA<<<g, b, 0, streamA>>>(d_x);
CUDA_CHECK(cudaEventRecord(e, streamA));            // stamps e at the end of A's work
CUDA_CHECK(cudaStreamWaitEvent(streamB, e, 0));     // B will wait for e before continuing
kernelB<<<g, b, 0, streamB>>>(d_x);                 // depends on A's result, on a different stream

שימו לב ל-cudaEventCreateWithFlags(&e, cudaEventDisableTiming): כשה-event משמש רק לסנכרון ולא למדידה, הדגל הזה מכבה את מנגנון החותמת ומוזיל את הoverhead. חשוב: cudaStreamWaitEvent חוסם רק את ה-GPU-side של streamB; הhost ממשיך מיד הלאה. זו בדיוק המהות - סנכרון בין זרמים בלי לערב את ה-CPU.

שלוש דרגות סנכרון של הhost. כשהhost כן צריך להמתין, יש שלוש רמות גרעיניות, מהצרה לרחבה:

קריאה ממתינה ל חוסמת את הhost?
cudaEventSynchronize(e) שה-event e הספציפי יוקלט כן, עד ל-e
cudaStreamSynchronize(s) שכל העבודה ב-stream s תסתיים כן, עד לריקון s
cudaDeviceSynchronize() שכל העבודה בכל ה-streams תסתיים כן, עד לריקון הdevice כולו

הכלל: השתמשו בקריאה הצרה ביותר שמספיקה. cudaDeviceSynchronize הוא הפטיש הכבד - הוא בולם את כל הקבילות ולכן שמור לסוף התוכנית או ל-debug. בצנרת קבילות מסנכרנים בדרך כלל ברמת ה-stream (cudaStreamSynchronize על כל זרם), או אפילו ברמת ה-event אם רוצים להמתין לאבן-דרך אחת בלבד. תוספת שימושית: cudaEventQuery(e) ו-cudaStreamQuery(s) בודקים בלי לחסום האם הסתיים - מחזירים cudaSuccess אם כן, cudaErrorNotReady אם עדיין רץ - ומאפשרים לhost לעשות עבודה מועילה בזמן שהוא "מציץ".

צנרת חופפת - a multi-stream pipeline

עכשיו נחבר הכל לתבנית שהיא הסיבה שלמדנו pinned, async, streams ו-events: צנרת שמחפפת H2D, חישוב ו-D2H על פני נתחים. במקום להעתיק מערך ענק כולו, לחשב, ולהחזיר - שלושה שלבים סדרתיים - נפרק אותו ל-chunks (נתחים), נשבץ אותם ב-streams מתחלפים, ובזמן שה-GPU מחשב על נתח אחד, ה-Copy Engine מעתיק את הבא:

const int nStreams = 4;
cudaStream_t st[nStreams];
for (int s = 0; s < nStreams; s++)
    CUDA_CHECK(cudaStreamCreateWithFlags(&st[s], cudaStreamNonBlocking));

int chunk = n / nChunks;                     // chunk size (assumes an even split)
size_t cb = (size_t)chunk * sizeof(float);

for (int c = 0; c < nChunks; c++) {
    int off = c * chunk;                     // chunk offset
    cudaStream_t s = st[c % nStreams];       // round-robin between the streams
    CUDA_CHECK(cudaMemcpyAsync(d_a+off, h_a+off, cb, cudaMemcpyHostToDevice, s));
    CUDA_CHECK(cudaMemcpyAsync(d_b+off, h_b+off, cb, cudaMemcpyHostToDevice, s));
    vecAdd<<<(chunk+255)/256, 256, 0, s>>>(d_a+off, d_b+off, d_c+off, chunk);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaMemcpyAsync(h_c+off, d_c+off, cb, cudaMemcpyDeviceToHost, s));
}
for (int s = 0; s < nStreams; s++)
    CUDA_CHECK(cudaStreamSynchronize(st[s]));    // wait for all streams to finish

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

   Serial (single stream, no overlap):
   H2D₀ ─► K₀ ─► D2H₀ ─► H2D₁ ─► K₁ ─► D2H₁ ─► ...
   |<------------ everything one after another ------------>|

   Overlapped (several streams, pinned):
   Copy H2D:  H2D₀  H2D₁  H2D₂  H2D₃
   Compute :        K₀    K₁    K₂    K₃
   Copy D2H:              D2H₀  D2H₁  D2H₂  D2H₃
              ▲ the three paths work in parallel - the H100 holds several Copy Engines

חשבון מודל. נסמן ב-t את זמן העתקת נתח (בכיוון אחד) ונניח שהוא שווה לזמן החישוב עליו, ונקח 4 נתחים. הגרסה הסדרתית לחלוטין - לכל נתח H2D ואז K ואז D2H, אחד אחרי השני - עולה בערך 3t x 4 = 12t. עם חפיפה מלאה של שלושת הנתיבים, הזמן מתקרב ל-max(סך ה-H2D, סך החישוב, סך ה-D2H) ועוד "זנב" של מילוי וריקון הצנרת בקצוות - בערך 4t + 2t = 6t. חיסכון של כ-50%. ככל שמרבים נתחים, הזנב נעשה זניח והזמן מתקרב לחסם התיאורטי max(copy, compute) - זמן הרכיב האיטי לבדו, כאילו שני האחרים חינם. שני תנאים הכרחיים שקל לשכוח: (1) h_a, h_b, h_c חייבים להיות pinned, אחרת ה-async מתנוון; (2) העבודה חייבת לשבת על streams לא-default (רצוי cudaStreamNonBlocking), אחרת ה-default stream מסדרן הכל. את שני המספרים - 12t הסדרתי מול ~6t החופף - תמדדו בעצמכם בתרגול עם events.

הערה לקראת 5.5: הצנרת הזאת מנפיקה, בכל איטרציה, עשרות launches והעתקות מהhost. כשמריצים אותה שוב ושוב בלולאה חמה (כמו ב-inference), overhead ההנפקה מהhost, מיקרו-שנייה לlaunch, מצטברת. זו בדיוק הבעיה ש-CUDA Graphs באים לפתור: לוכדים את כל רצף העבודה על stream פעם אחת ומשגרים אותו כיחידה אחת. נבנה את זה בשיעור 5.5.

סיכום

  • כמעט כל קריאת Runtime מחזירה cudaError_t; cudaGetErrorString/cudaGetErrorName מפענחים אותה, ו-cudaPeekAtLastError (מציץ) מול cudaGetLastError (מציץ ומאפס) ניגשים לשגיאה האחרונה per-thread.
  • לlaunch kernel שני סוגי שגיאות: שגיאת launch (תצורה) שנתפסת מיד ב-cudaGetLastError(), ושגיאת הרצה (out-of-bounds) שצצה רק בקריאה המסנכרנת הבאה; לכן בדיקה מלאה = cudaGetLastError() ואחריו סנכרון.
  • שגיאה דביקה (sticky), הנובעת מ-kernel שהשחית את ה-context, מזהמת כל קריאה עוקבת ולא ניתנת לניקוי; ההתאוששות היחידה היא הרג התהליך או cudaDeviceReset().
  • ה-Runtime מאותחל עצלנית בקריאה הראשונה, וטעינת מודולים עצלה (CUDA_MODULE_LOADING=LAZY) דוחה את טעינת ה-SASS לlaunch הראשונה; לכן תמיד מריצים איטרציית חימום לפני מדידה.
  • מאקרו CUDA_CHECK יחיד שעוטף כל קריאה, מדפיס קובץ, שורה, הקריאה (#call) וההודעה ואז יוצא, הוא תנאי הכרחי לקוד CUDA שאפשר לתחזק; העלות זניחה והתועלת עצומה.
  • בתוך stream הפעולות מסודרות, בין streams אין סדר - וזה המנוף לקבילות; ה-legacy default stream מסנכרן את כל הזרמים החוסמים, בעוד per-thread default stream (nvcc --default-stream per-thread) נותן לכל host thread זרם ברירת-מחדל בלתי-חוסם משלו.
  • cudaMemcpyAsync על stream חוזרת מיד, אך רק מזיכרון pinned וב-stream לא-default היא באמת אסינכרונית; אחרת היא מתנוונת בשקט ל-copy חוסם.
  • cudaEvent עם cudaEventRecord ו-cudaEventElapsedTime מודד זמן GPU ברזולוציה תת-מיקרו-שנייה (זכרו להעביר את ה-stream ל-record ולהמתין ב-cudaEventSynchronize לפני קריאת ההפרש).
  • cudaStreamWaitEvent בונה תלות צולבת בין streams בלי לחסום את הhost, ולסנכרון הhost יש שלוש דרגות - cudaEventSynchronize (event), cudaStreamSynchronize (זרם), cudaDeviceSynchronize (device כולו) - שמהן בוחרים את הצרה ביותר.
  • צנרת חופפת על פני נתחים וכמה streams מקרבת את הזמן מ-H2D + compute + D2H הסדרתי אל max(copy, compute), ודורשת בו-זמנית pinned, async ו-streams לא-default; זו גם התשתית ל-CUDA Graphs בשיעור 5.5.