5.3 Runtime API, streams ו events פתרון
פתרון - Runtime API, streams ו-events¶
הערה: הזמנים המוחלטים, רוחבי הפס ויחסי ה-speedup בפלטים למטה הם דוגמאות ממכונת H100 SXM עם PCIe Gen5, וישתנו אצלכם לפי הכרטיס, גרסת ה-PCIe ומצב המערכת. מה שלא ישתנה הוא המבנה, סדרי הגודל, ויחסי ה-speedup בין הגרסאות - עליהם התרגיל בנוי. כל הקטעים מתקמפלים עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4, sm_80 על A100). בכל הקבצים מניחים שהמאקרו CUDA_CHECK וה-kernel vecAdd מ-הכנה נמצאים בראש הקובץ.
פתרון תרגיל 1 - מאקרו CUDA_CHECK ולכידת שגיאה אסינכרונית¶
int main(void) {
int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc(&d_a, bytes));
CUDA_CHECK(cudaMalloc(&d_b, bytes));
CUDA_CHECK(cudaMalloc(&d_c, bytes));
// (2) valid launch
vecAdd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaGetLastError()); // launch error (none)
CUDA_CHECK(cudaDeviceSynchronize()); // execution error (none)
printf("[ok] valid launch passed\n");
// (3) launch error: 1025 threads per block (over 1024)
vecAdd<<<1, 1025>>>(d_a, d_b, d_c, n);
cudaError_t e1 = cudaGetLastError(); // caught immediately
printf("[launch] %s\n", cudaGetErrorName(e1));
// (4) execution error: a fake n that's 4x too large, and the grid is also enlarged by that same fake n
// so that millions of threads truly index far beyond the buffers of 2^20 elements
vecAdd<<<(n * 4 + 255) / 256, 256>>>(d_a, d_b, d_c, n * 4);
cudaError_t e2 = cudaGetLastError(); // still clean here!
printf("[after launch] %s\n", cudaGetErrorName(e2));
cudaError_t e3 = cudaDeviceSynchronize(); // only here does it surface
printf("[after sync] %s\n", cudaGetErrorName(e3));
// (5) sticky error: an unrelated subsequent call fails too
float* extra;
cudaError_t e4 = cudaMalloc(&extra, bytes);
printf("[sticky] cudaMalloc -> %s\n", cudaGetErrorName(e4));
return 0;
}
הcompilation והרצה:
פלט צפוי:
[ok] valid launch passed
[launch] cudaErrorInvalidConfiguration
[after launch] cudaSuccess
[after sync] cudaErrorIllegalAddress
[sticky] cudaMalloc -> cudaErrorIllegalAddress
שימו לב לשורה השלישית והרביעית: מיד אחרי הlaunch הפגומה cudaGetLastError החזיר cudaSuccess - שגיאת ההרצה עוד לא צפה. רק cudaDeviceSynchronize חשף את cudaErrorIllegalAddress. השורה החמישית מוכיחה את הדביקות: cudaMalloc תמים לגמרי, שאין לו קשר לגישה מחוץ לתחום, מחזיר את אותה שגיאה בדיוק, כי ה-context הושחת.
למה זה עבד: שגיאת תצורה (1025 > 1024) נבדקת על ידי הhost בעת הlaunch, ולכן היא סינכרונית ונתפסת ב-cudaGetLastError מיד. גישה מחוץ לתחום קורית על ה-GPU בזמן הריצה, זמן רב אחרי שהhost המשיך, ולכן היא אסינכרונית וצפה רק בנקודת הסנכרון הבאה - ומשם היא דביקה ומזהמת כל קריאה.
איך להכליל: השתמשו תמיד בזוג cudaGetLastError() (מיד אחרי כל launch) ו-cudaDeviceSynchronize() (ב-build של debug) כדי לתפוס את שני סוגי השגיאות. ברגע שאתם רואים שגיאה דביקה מתגלגלת לקריאה אקראית, דעו שהמקור האמיתי הוא kernel קודם - הריצו תחת compute-sanitizer ./streams כדי לאתר את הגישה החורגת המדויקת.
פתרון תרגיל 2 - מדידת זמן kernel מדויקת עם events¶
int main(void) {
int n = 1 << 26; // 67,108,864 elements
size_t bytes = (size_t)n * sizeof(float); // 256 MiB
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc(&d_a, bytes));
CUDA_CHECK(cudaMalloc(&d_b, bytes));
CUDA_CHECK(cudaMalloc(&d_c, bytes));
CUDA_CHECK(cudaMemset(d_a, 0, bytes));
CUDA_CHECK(cudaMemset(d_b, 0, bytes));
int threads = 256, blocks = (n + threads - 1) / threads;
// warm-up
vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
float best = 1e30f;
for (int r = 0; r < 20; r++) {
CUDA_CHECK(cudaEventRecord(start));
vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
if (ms < best) best = ms;
}
double gb = (double)n * 12.0 / 1e9; // 3 arrays x 4 bytes
printf("vecAdd: %.3f ms -> %.0f GB/s effective\n", best, gb / (best / 1e3));
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
return 0;
}
פלט צפוי:
החישוב ידנית: n * 12 = 67,108,864 x 12 = 805,306,368 בתים ≈ 0.805 GB; חלקי 1.203/1000 שניות ≈ 669 GB/s. מול 3.35 TB/s של ה-HBM3, זה כ-20%.
למה זה עבד: ה-event מוקלט על ה-GPU בין הפעולות באותו stream, ולכן cudaEventElapsedTime מודד את זמן ה-GPU נטו, בלי overhead הhost, ברזולוציה של כחצי מיקרו-שנייה. איטרציית החימום ספגה את ה-lazy init ואת ה-JIT, והמינימום על פני 20 חזרות ניקה רעש מערכת. רוחב הפס הנמוך יחסית (20% מהשיא) מאשר ש-vecAdd memory-bound: ה-arithmetic intensity שלו זעומה (חיבור אחד לכל 12 בתים), אז החומרה מוגבלת ברוחב הפס ולא בחישוב.
איך להכליל: לכל kernel, מדדו זמן עם events, גזרו ממנו את רוחב הפס האפקטיבי (בתים שנגעתם בהם חלקי זמן) או את ה-FLOP/s (פעולות חלקי זמן), והשוו לשיא של הכרטיס כדי לדעת אם אתם memory-bound או compute-bound - זו נקודת ההתחלה של כל אופטימיזציה, ונחזור אליה עם מודל ה-Roofline בפרק 8.
פתרון תרגיל 3 - צנרת חופפת מול סדרתית ומדידת speedup¶
static float timeSerial(float* hc, float* ha, float* hb,
float* da, float* db, float* dc, int n, size_t bytes) {
int threads = 256, blocks = (n + threads - 1) / threads;
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaEventRecord(s));
CUDA_CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
vecAdd<<<blocks, threads>>>(da, db, dc, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return ms;
}
static float timeOverlap(float* hc, float* ha, float* hb,
float* da, float* db, float* dc, int n, size_t bytes,
int nStreams, int nChunks) {
cudaStream_t st[8];
for (int i = 0; i < nStreams; i++)
CUDA_CHECK(cudaStreamCreateWithFlags(&st[i], cudaStreamNonBlocking));
int chunk = n / nChunks;
size_t cb = (size_t)chunk * sizeof(float);
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaEventRecord(s));
for (int c = 0; c < nChunks; c++) {
int off = c * chunk;
cudaStream_t str = st[c % nStreams];
CUDA_CHECK(cudaMemcpyAsync(da+off, ha+off, cb, cudaMemcpyHostToDevice, str));
CUDA_CHECK(cudaMemcpyAsync(db+off, hb+off, cb, cudaMemcpyHostToDevice, str));
vecAdd<<<(chunk+255)/256, 256, 0, str>>>(da+off, db+off, dc+off, chunk);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpyAsync(hc+off, dc+off, cb, cudaMemcpyDeviceToHost, str));
}
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
for (int i = 0; i < nStreams; i++) CUDA_CHECK(cudaStreamDestroy(st[i]));
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return ms;
}
int main(void) {
int n = 1 << 26;
size_t bytes = (size_t)n * sizeof(float);
float *ha, *hb, *hc, *ref;
CUDA_CHECK(cudaMallocHost(&ha, bytes)); // pinned
CUDA_CHECK(cudaMallocHost(&hb, bytes));
CUDA_CHECK(cudaMallocHost(&hc, bytes));
ref = (float*)malloc(bytes);
for (int i = 0; i < n; i++) { ha[i] = 1.0f; hb[i] = 2.0f; ref[i] = 3.0f; }
float *da, *db, *dc;
CUDA_CHECK(cudaMalloc(&da, bytes));
CUDA_CHECK(cudaMalloc(&db, bytes));
CUDA_CHECK(cudaMalloc(&dc, bytes));
float ts = timeSerial(hc, ha, hb, da, db, dc, n, bytes);
printf("serial : %.2f ms\n", ts);
for (int ns = 2; ns <= 4; ns++) {
float to = timeOverlap(hc, ha, hb, da, db, dc, n, bytes, ns, 8);
// correctness check
int bad = 0;
for (int i = 0; i < n; i++) if (fabsf(hc[i] - ref[i]) > 1e-5f) { bad = 1; break; }
printf("overlap ns=%d: %.2f ms speedup %.2fx %s\n",
ns, to, ts / to, bad ? "MISMATCH" : "ok");
}
return 0;
}
פלט צפוי:
serial : 15.87 ms
overlap ns=2: 11.03 ms speedup 1.44x ok
overlap ns=3: 10.42 ms speedup 1.52x ok
overlap ns=4: 10.31 ms speedup 1.54x ok
הסדרתי מורכב מ-H2D של שני מערכים (~9.8 ms), kernel (~1.2 ms) ו-D2H (~4.9 ms) בזה אחר זה. בגרסה החופפת, ה-H2D של נתח אחד רץ בעוד ה-kernel של הנתח הקודם מחשב וה-D2H של נתח מוקדם עוד יותר יורד - שלושת הנתיבים במקביל, והזמן מתקרב לרכיב הכבד ביותר (ה-H2D) לבדו.
ניסוי המלכוד (סעיף 6) עם malloc pageable:
ה-speedup התאדה: עם זיכרון pageable, cudaMemcpyAsync נופלת חזרה ל-bounce buffer ומסתנכרנת, כך שאין שום חפיפה למרות שהזרמים קיימים.
למה זה עבד: החפיפה נשענת על שלושה תנאים בו-זמנית - זיכרון pinned (מאפשר DMA ישיר ואמיתית-אסינכרוני), cudaMemcpyAsync (מנפיק בלי לחסום), ו-streams לא-default (מאפשרים קבילות). ה-H100 מחזיק כמה Copy Engines עצמאיים מה-SMs, ולכן הוא מבצע H2D, חישוב ו-D2H פיזית במקביל. ה-speedup מתייצב סביב פי 1.5 כי vecAdd memory-bound וה-PCIe הוא צוואר הבקבוק - חפיפת ההעתקות עם החישוב הזעיר חוסכת בעיקר את זמן ה-kernel שנבלע.
איך להכליל: כל workload שאפשר לחלק לנתחים בלתי-תלויים נהנה מהתבנית הזאת. הרווח הגדול ביותר מתקבל כשזמני ה-copy וה-compute דומים; אם החישוב כבד בהרבה מההעתקה, ההעתקה כבר "מתחבאת" וה-speedup קטן, ואם ההעתקה שולטת - כמו כאן - החפיפה חוסכת את זמן החישוב. מספר הנתחים צריך להיות גדול דיו כדי למלא את הצנרת אך לא כה גדול שoverhead ההנפקה מהhost תשתלט.
פתרון תרגיל 4 - תלות צולבת בין זרמים עם cudaStreamWaitEvent¶
__global__ void scaleK(float* x, int n) { int i = blockIdx.x*blockDim.x+threadIdx.x; if (i<n) x[i] *= 2.0f; }
__global__ void addOneK(float* x, int n) { int i = blockIdx.x*blockDim.x+threadIdx.x; if (i<n) x[i] += 1.0f; }
int main(void) {
int n = 1 << 22;
size_t bytes = (size_t)n * sizeof(float);
float* hx; CUDA_CHECK(cudaMallocHost(&hx, bytes));
for (int i = 0; i < n; i++) hx[i] = 5.0f; // x_init = 5
float* dx; CUDA_CHECK(cudaMalloc(&dx, bytes));
cudaStream_t sA, sB;
CUDA_CHECK(cudaStreamCreateWithFlags(&sA, cudaStreamNonBlocking));
CUDA_CHECK(cudaStreamCreateWithFlags(&sB, cudaStreamNonBlocking));
cudaEvent_t e;
CUDA_CHECK(cudaEventCreateWithFlags(&e, cudaEventDisableTiming));
int th = 256, bl = (n + th - 1) / th;
CUDA_CHECK(cudaMemcpyAsync(dx, hx, bytes, cudaMemcpyHostToDevice, sA));
scaleK<<<bl, th, 0, sA>>>(dx, n); // sA: x *= 2 -> 10
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaEventRecord(e, sA)); // stamps e at the end of scale
CUDA_CHECK(cudaStreamWaitEvent(sB, e, 0)); // sB will wait for e
addOneK<<<bl, th, 0, sB>>>(dx, n); // sB: x += 1 -> 11
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpyAsync(hx, dx, bytes, cudaMemcpyDeviceToHost, sB));
CUDA_CHECK(cudaStreamSynchronize(sB));
printf("x[0] = %.1f (expected 11.0)\n", hx[0]);
return 0;
}
פלט צפוי:
התוצאה 11 = 5*2 + 1 מוכיחה שה-addOneK ב-sB רץ אחרי ה-scaleK ב-sA. אם תסירו את שורת cudaStreamWaitEvent, ריצות חוזרות עלולות להחזיר 11 לפעמים ו-6 לפעמים (5 + 1, כש-addOne רץ לפני scale) - תוצאה לא-דטרמיניסטית.
למה זה עבד: cudaEventRecord(e, sA) מניח את e בזנב תור העבודה של sA, ו-e נחתם רק כשה-scaleK מסתיים. cudaStreamWaitEvent(sB, e, 0) מכניס לתור של sB מחסום שממתין לחתימת e, ולכן ה-addOneK לא מתחיל לפני שה-scale גמר - למרות ששני ה-kernels בזרמים שונים. הhost, לעומת זאת, לא נחסם כלל בשלב הזה; רק ה-GPU-side של sB ממתין.
איך להכליל: cudaStreamWaitEvent הוא הכלי לבנות גרף תלויות בין זרמים - למשל, מספר יצרנים בזרמים שונים שכולם צריכים להסתיים לפני צרכן יחיד, או פיצול-ומיזוג (fork-join) של עבודה מקבילית - בלי לשרשר הכל לזרם אחד (שהורג קבילות) ובלי להעיר את ה-CPU (שמוסיף latency). זו בדיוק הסמנטיקה שעליה בנויים CUDA Graphs בשיעור 5.5: גרף התלויות שכאן בונים ידנית, שם נלכד ומשוגר כיחידה אחת.
פתרון תרגיל 5 (בונוס) - פרופיל של הצנרת ב-Nsight Systems¶
קטע מייצג מסיכום ה-stats (הערכים משתנים):
** CUDA GPU MemOps Summary (by Time) (cuda_gpu_mem_time_sum):
Time(%) Total Time Count Operation
------- ---------- ----- ---------------------------
66.4 9,812,000 16 [CUDA memcpy Host-to-Device]
33.1 4,905,000 8 [CUDA memcpy Device-to-Host]
** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):
Time(%) Total Time Count Name
------- ---------- ----- --------
100.0 1,190,000 8 vecAdd(...)
בציר הזמן של ה-GUI תראו שלוש שורות פעילות בו-זמנית: שורת ה-HtoD, שורת ה-kernels, ושורת ה-DtoH, כשהפסים חופפים בזמן (נתח נמצא ב-D2H בעוד נתח אחר ב-kernel ושלישי ב-H2D). בגרסה הסדרתית, לעומת זאת, כל הפסים צמודים ורצופים בשורה אחת, בלי חפיפה.
למה זה עבד: nsys דוגם את ה-CUDA API ואת פעילות ה-GPU ומצייר אותם על ציר זמן משותף לפי stream, כך שחפיפה נראית ישירות כפסים מקבילים בזרמים שונים. סיכום ה-stats מאשר שהעבודה אכן פוזרה (8 kernels, 16 העתקות H2D לשני מערכים, 8 העתקות D2H) - נתח לכל stream ב-round-robin.
איך להכליל: כשה-speedup שמדדתם עם events נמוך מהצפוי, nsys הוא הכלי הראשון שאליו פונים - הוא מראה למה אין חפיפה (זרם default שמסדרן, זיכרון pageable שמסנכרן, או תלות סמויה בין הפעולות). את הפרופיילינג המעמיק, ברמת ה-kernel הבודד עם ncu, נפרוש בפרק 9.