לדלג לתוכן

5.3 Runtime API, streams ו events תרגול

תרגול - Runtime API, streams ו-events

בתרגול הזה תבנו את שלושת העמודים של תכנות Runtime API מקצועי: טיפול חסין בשגיאות, מדידת זמן מדויקת עם events, וקבילות אמיתית עם streams. תעברו מקוד שנכשל בשקט לקוד שצועק בדיוק היכן נשבר; תמדדו kernel ברזולוציה תת-מיקרו-שנייה; תבנו צנרת שמחפפת H2D, חישוב ו-D2H ותמדדו את ה-speedup מול הגרסה הסדרתית; ולבסוף תבנו תלות צולבת בין שני זרמים עם cudaStreamWaitEvent. עבדו לפי הסדר - כל תרגיל נשען על הקודם. כל התרגילים רצים על GPU של NVIDIA; אם אין לכם כרטיס מקומי, השתמשו ב-Google Colab (T4 חינם) והחליפו את דגל הארכיטקטורה בהתאם.

הכנה

צרו קובץ עבודה streams.cu. בראשו שימו את המאקרו והכותרות שבהם נשתמש בכל התרגילים:

#include <cstdio>
#include <cstdlib>
#include <cmath>
#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)

__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];
}

הcompilation (על H100; החליפו ל-sm_75 על T4, sm_80 על A100):

nvcc -O2 -arch=sm_90a streams.cu -o streams

תרגיל 1 - מאקרו CUDA_CHECK ולכידת שגיאה אסינכרונית

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

  1. כתבו main שמקצה d_a, d_b, d_c בגודל n = 1<<20 עם cudaMalloc, עוטף כל קריאה ב-CUDA_CHECK.
  2. השיקו את vecAdd תקין עם תצורה חוקית, ואז הריצו CUDA_CHECK(cudaGetLastError()) ו-CUDA_CHECK(cudaDeviceSynchronize()). ודאו שהתוכנית רצה נקי.
  3. עכשיו זרעו שגיאת launch: השיקו kernel עם 1025 threads לבלוק (מעל המקסימום 1024). הוסיפו cudaGetLastError() מיד אחרי הlaunch ובדקו איזו שגיאה נתפסת.
  4. החזירו את התצורה לחוקית, אך זרעו שגיאת הרצה: השיקו vecAdd עם n גדול מגודל ההקצאה האמיתית (למשל העבירו ל-kernel n כפול 4), כך שייגש מחוץ לתחום. שימו לב היכן בדיוק השגיאה נתפסת - האם מיד אחרי הlaunch, או רק ב-cudaDeviceSynchronize?
  5. אחרי שנתפסה שגיאת ההרצה, נסו להוסיף עוד cudaMalloc אחריה (בלי exit) ובדקו: האם גם היא נכשלת? הסבירו במונחי "שגיאה דביקה".

רמז: שגיאת launch (תצורה) נתפסת ב-cudaGetLastError() מיד; שגיאת הרצה (out-of-bounds) צצה רק בקריאה המסנכרנת הבאה. כדי לראות את סעיף 5 בלי ש-CUDA_CHECK יסיים את התוכנית, בדקו שם ידנית עם cudaError_t e = cudaMalloc(...) והדפיסו את cudaGetErrorName(e).

תרגיל 2 - מדידת זמן kernel מדויקת עם events

השתמשו ב-cudaEvent כדי למדוד את vecAdd וחשבו ממנו רוחב פס אפקטיבי.

  1. הקצו והאתחלו d_a, d_b, d_c בגודל n = 1<<26 (כ-67 מיליון איברים). מלאו את a,b בערכים כלשהם על ה-GPU או העתיקו מהhost.
  2. הריצו איטרציית חימום אחת של vecAdd ואחריה cudaDeviceSynchronize, בלי למדוד אותה.
  3. צרו start ו-stop עם cudaEventCreate. הקליטו start, השיקו את vecAdd, הקליטו stop, סנכרנו את stop, וקראו cudaEventElapsedTime.
  4. הריצו את המדידה בלולאה של 20 חזרות ושמרו את הזמן המינימלי (הכי פחות רעש). הדפיסו אותו במילישניות.
  5. חשבו את רוחב הפס האפקטיבי: vecAdd נוגע ב-3 מערכי float (2 קריאות, כתיבה אחת), כלומר n * 12 בתים. חלקו בזמן (בשניות) והדפיסו ב-GB/s. השוו לרוחב הפס של ה-HBM3 ב-H100 (בסדר גודל של 3.35 TB/s) - איזה אחוז יצא, והאם ה-kernel memory-bound?

רמז: cudaEventElapsedTime(&ms, start, stop) מחזיר float מילישניות ודורש ש-stop כבר הוקלט - מכאן cudaEventSynchronize(stop) לפניו. אל תשכחו לבדוק cudaGetLastError() אחרי הlaunch. הזמן המינימלי על פני חזרות מנקה רעש מערכת.

תרגיל 3 - צנרת חופפת מול סדרתית ומדידת speedup

זהו לב התרגול: תבנו שתי גרסאות של אותה עבודה ותמדדו כמה מהיר החופף.

  1. הקצו זיכרון host pinned ל-h_a, h_b, h_c עם cudaMallocHost, בגודל n = 1<<26. אתחלו a,b.
  2. גרסה סדרתית (baseline): cudaMemcpy H2D של כל a ו-b, launch אחת של vecAdd על כל n, cudaMemcpy D2H של c. עטפו את כל הבלוק ב-events ומדדו את זמנו הכולל.
  3. גרסה חופפת: פרקו את n ל-nChunks (התחילו ב-8) ופזרו round-robin על nStreams זרמים שנוצרו עם cudaStreamCreateWithFlags(..., cudaStreamNonBlocking). לכל נתח: cudaMemcpyAsync H2D של a,b, launch vecAdd על הנתח, cudaMemcpyAsync D2H של c - הכל ב-stream של אותו נתח. מדדו את זמן כל הצנרת עם events (הקליטו start לפני הלולאה ו-stop אחריה, וסנכרנו).
  4. אמתו נכונות: השוו את h_c מהגרסה החופפת מול הגרסה הסדרתית (או מול חישוב CPU), ודאו שהם זהים.
  5. הדפיסו את זמן הסדרתי, זמן החופף, ואת ה-speedup = t_serial / t_overlap. הריצו עם nStreams של 2, 3, ו-4 ורשמו כיצד ה-speedup משתנה.
  6. ניסוי המלכוד: החליפו את cudaMallocHost ל-malloc רגיל (pageable) בגרסה החופפת בלבד, והריצו שוב. מה קרה ל-speedup, ולמה?

רמז: החפיפה דורשת בו-זמנית שלושה תנאים - זיכרון pinned, cudaMemcpyAsync, ו-streams לא-default. אם אחד מהם חסר, אין חפיפה. ה-speedup הצפוי כשזמן ההעתקה וזמן החישוב דומים הוא סביב פי 1.5 עד פי 2 - vecAdd הוא memory-bound, כך שה-H2D וה-D2H (על ה-PCIe) הם הרכיב הכבד, וחפיפתם עם החישוב היא הרווח.

תרגיל 4 - תלות צולבת בין זרמים עם cudaStreamWaitEvent

בנו תלות מפורשת בין שני זרמים בלי לחסום את הhost.

  1. צרו שני זרמים sA ו-sB (cudaStreamNonBlocking) ו-event יחיד e עם הדגל cudaEventDisableTiming.
  2. ב-sA השיקו kernel scale שמכפיל מערך d_x בקבוע (כתבו kernel פשוט x[i] *= 2.0f).
  3. הקליטו את ה-event: cudaEventRecord(e, sA).
  4. הורו ל-sB להמתין: cudaStreamWaitEvent(sB, e, 0), ואז ב-sB השיקו kernel addOne שמוסיף 1 לכל איבר של אותו d_x.
  5. סנכרנו, העתיקו d_x לhost ובדקו שהתוצאה היא x_init * 2 + 1 לכל איבר - כלומר ש-sB באמת חיכה ל-sA.
  6. כדי לוודא שהתלות נחוצה, הסירו זמנית את cudaStreamWaitEvent והריצו כמה פעמים; האם התוצאה יכולה להשתבש? הסבירו מדוע (מרוץ נתונים בין הזרמים).

רמז: cudaStreamWaitEvent(sB, e, 0) חוסם רק את צד ה-GPU של sB עד שה-event נחתם ב-sA; הhost לא נחסם. בלי ה-wait, addOne ב-sB יכול לרוץ לפני, אחרי, או במקביל ל-scale ב-sA - התוצאה תלויה בתזמון, וזו הגדרת מרוץ נתונים.

תרגיל 5 (בונוס) - פרופיל של הצנרת ב-Nsight Systems

אמתו חזותית שהחפיפה באמת קורית.

  1. הריצו את גרסת הצנרת החופפת מתרגיל 3 תחת nsys: nsys profile -o pipe --stats=true ./streams.
  2. פתחו את pipe.nsys-rep ב-Nsight Systems GUI (או קראו את סיכום ה-stats בטרמינל), ואתרו את שורות ה-streams.
  3. חפשו בציר הזמן את החפיפה: פסי ה-Memcpy HtoD, ה-kernels, וה-Memcpy DtoH צריכים להיות מדורגים - כלומר לרוץ בו-זמנית בזרמים שונים - ולא בשורה אחת רציפה.
  4. השוו לגרסה הסדרתית: פרפלו גם אותה וראו שהפסים שם צמודים ורציפים, בלי חפיפה.

רמז: ב-nsys stats חפשו את הטבלאות cuda_gpu_kern_sum ו-cuda_gpu_mem_time_sum; ב-GUI, שורת כל stream נפרדת בציר הזמן, וחפיפה נראית כפסים שמתחילים לפני שהקודם הסתיים. אם אין חפיפה גם בגרסה ה"חופפת", חזרו לבדוק שהזיכרון pinned ושהזרמים אינם default.