5.5 CUDA Graphs לצמצום overhead תרגול

תרגול - CUDA Graphs לצמצום overhead

בתרגול הזה תרגישו במו ידיכם את overhead הlaunch ואת הריפוי שלה בעזרת גרף CUDA. תתחילו בבניית לולאה שמשגרת מאות kernels זעירים ותמדדו את הזמן הכולל ואת הoverhead לכל launch, ואז תלכדו את אותו רצף בדיוק לתוך גרף ותמדדו כמה ירדה הoverhead - לרצפים של הרבה launches קטנות תראו נפילה גדולה. אחר כך תחקרו את המגבלה החדה של הגרף (המצביעים הצרובים) על-ידי שחרור והקצאה מחדש של ה-buffer, ולבסוף תסקיצו איך torch.cuda.graph היה חל על צעד inference. עבדו לפי הסדר - כל תרגיל בונה על קודמו. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA (רק המספרים המוחלטים ישתנו).

הכנה

צרו קובץ graphs.cu. לאורך כל התרגול נשתמש ב-kernel זעיר במכוון, כדי שoverhead הlaunch תשלוט בזמן ולא החישוב. הקומפילציה (החליפו את sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):

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

מקרו הבדיקה וה-kernel הזעיר שנשתמש בהם:

#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  '%s'\n  -> %s\n",         \
                    __FILE__, __LINE__, #call, cudaGetErrorString(err_));     \
            exit(EXIT_FAILURE);                                               \
        }                                                                     \
    } while (0)

__global__ void tiny(float* d) { d[0] += 1.0f; }   // negligible work - overhead dominates

תרגיל 1 - מדידת overhead הlaunch בלולאה

1. הקצו float* d בודד על הdevice עם cudaMalloc ואפסו אותו עם cudaMemset.
2. צרו stream לא-default עם cudaStreamCreate.
3. כתבו לולאה שמשגרת את tiny<<<1,1,0,stream>>>(d) בדיוק N = 512 פעמים, ואחריה cudaStreamSynchronize. עטפו את כל הלולאה (512 הlaunches) במדידת cudaEvent.
4. הריצו איטרציית חימום אחת של כל הלולאה לפני המדידה, ואז מדדו את החציון (או המינימום) על פני 50 חזרות.
5. חשבו ודווחו את הoverhead לכל launch: זמן_כולל_במיקרו / 512.

רמז: השתמשו ב-stream לא-default (לא stream 0), אחרת בהמשך לא תוכלו ללכוד. את זמן ה-cudaEvent מחלקים ב-512 כדי לקבל את הoverhead לlaunch בודדת; אתם אמורים לראות מספר בסדר גודל של יחידות מיקרו-שניות (למשל 3-10 us). זכרו שה-kernel עצמו כמעט לא עולה זמן ריצה - כמעט כל הזמן שתמדדו הוא host overhead.

תרגיל 2 - לכידת אותו רצף לגרף ומדידת הריפוי

1. לכדו בדיוק את אותו רצף של 512 launches tiny לתוך גרף: cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal), הלולאה של 512 הlaunches על ה-stream, ואז cudaStreamEndCapture(stream, &graph).
2. הגישו למכונה עם cudaGraphInstantiate(&graphExec, graph, 0).
3. מדדו עם cudaEvent את הזמן של cudaGraphLaunch(graphExec, stream) בודד (עם cudaStreamSynchronize אחריו), אחרי חימום, חציון על 50 חזרות.
4. דווחו את ה-speedup: זמן_הלולאה_מתרגיל_1 / זמן_ה-replay. אמתו נכונות: הריצו replay מספר ידוע של פעמים ובדקו ש-d[0] שווה למכפלה הצפויה.
5. שחררו הכל: cudaGraphExecDestroy, cudaGraphDestroy, cudaStreamDestroy, cudaFree.

רמז: בתוך הלכידה אסור להשתמש בשום פעולה חוסמת (לא cudaMemcpy רגילה, לא cudaDeviceSynchronize, לא cudaMalloc) - היא תפסול את הלכידה. ה-replay הבודד אמור להיות מהיר בהרבה מהלולאה, כי כל 512 הlaunches מוגשות בקריאת host אחת. לרצף של כל-כך הרבה kernels קטנים, ה-speedup יהיה גדול (בסדר גודל של פי כמה).

תרגיל 3 - סריקת מספר הlaunches

1. הפכו את N (מספר ה-kernels ברצף) לפרמטר, וחזרו על המדידה של תרגיל 1 (לולאה) ותרגיל 2 (גרף) עבור N ∈ {8, 64, 512, 4096}.
2. לכל N, דווחו שלוש עמודות: זמן הלולאה, זמן ה-replay, וה-speedup.
3. שרטטו (או תארו במילים) איך ה-speedup גדל עם N. הסבירו למה: מה בזמן הלולאה תלוי ב-N, ומה בזמן ה-replay תלוי ב-N?

רמז: זמן הלולאה גדל לינארית ב-N (כל launch מוסיפה את הoverhead שלה), בעוד שזמן ה-replay גדל הרבה יותר לאט - עלות ההגשה של הhost קבועה, ומה שגדל הוא רק זמן ההרצה של ה-kernels על הdevice. לכן ה-speedup עולה עם N. עבור N קטן (למשל 8) ייתכן שהגרף כמעט לא יעזור - הoverhead ממילא זניחה שם.

תרגיל 4 - המגבלה: מצביעים צרובים

1. קחו את הגרף בר-ההרצה מתרגיל 2 (שלכד launches על d). ודאו ש-replay אחד עובד ומעדכן את d[0].
2. עכשיו שחררו את d עם cudaFree(d), והקצו buffer חדש cudaMalloc(&d2, sizeof(float)). אפסו אותו.
3. הריצו שוב cudaGraphLaunch(graphExec, stream) והעתיקו את d2[0] לhost. האם d2 התעדכן? הסבירו את מה שקרה (או לא קרה).
4. אמתו את ההסבר שלכם: הקצו את d2 דרך שרשרת של הקצאות ובדקו עם compute-sanitizer ./graphs אם ה-replay נוגע בכתובת המשוחררת. תעדו את מה ש-compute-sanitizer מדווח.
5. הסבירו כיצד cudaGraphExecUpdate (או לכידה מחדש) היו הדרך הנכונה לתמוך ב-buffer חדש.

רמז: הכתובת של d נצרבה לתוך הגרף בזמן הלכידה. אחרי cudaFree(d), ה-replay ממשיך לכתוב לאותה כתובת ישנה - שכעת משוחררת. d2 לא יתעדכן (הוא בכתובת אחרת), וייתכן שתראו נזק שקט. הריצו compute-sanitizer ./graphs (הכלי מגיע עם ה-CUDA Toolkit) כדי לחשוף את ה-invalid access. זו הדגמה מוחשית לכלל "יציבות ההקצאות".

תרגיל 5 - השוואה ל-cudaEvent על ה-CPU מול ה-GPU

1. מדדו את הלולאה מתרגיל 1 בשתי דרכים במקביל: פעם עם cudaEvent (זמן GPU), ופעם עם שעון ה-CPU (std::chrono::steady_clock) שעוטף את הלולאה בלי cudaStreamSynchronize באמצע.
2. הסבירו למה מדידת ה-CPU של לולאת ההנפקה (בלי סנכרון) יכולה להיות קטנה מזמן ה-GPU: הhost מנפיק את הlaunches ורץ הלאה בזמן שה-GPU עדיין עובד.
3. עכשיו הוסיפו cudaStreamSynchronize בסוף ומדדו שוב עם ה-CPU. איך זה משתנה?
4. הסיקו מה מודד כל כלי: זמן ההנפקה של הhost מול זמן ההרצה בפועל על הdevice.

רמז: בלי סנכרון, שעון ה-CPU מודד רק כמה זמן לקח לhost לדחוף את 512 הlaunches לתור - זו בדיוק "overhead ההנפקה". עם סנכרון, הוא מודד גם את ההמתנה לסיום ההרצה. הפער בין השניים הוא כמה עבודה ה-GPU עשה במקביל להנפקה. cudaEvent תמיד מודד את זמן ה-GPU עצמו.

תרגיל 6 (בונוס) - סקיצת torch.cuda.graph לצעד inference

1. כתבו (בפייתון, כסקיצה - לא חובה להריץ מודל אמיתי) את דפוס ה-torch.cuda.graph לצעד inference: buffers סטטיים, חימום על stream צדדי, לכידה, ו-replay חוזר.
2. סמנו במפורש בקוד: היכן נצרבות הכתובות, מדוע static_input חייב להיות buffer קבוע, ולמה כותבים אליו עם copy_ במקום ליצור tensor חדש בכל צעד.
3. הסבירו במילים למה צעד decode של LLM (אותו forward pass של השכבות, שוב ושוב לכל טוקן) הוא מקרה אידאלי לגרף, וקשרו זאת לחשבון ה-480 launches לטוקן מההרצאה.
4. הסבירו מתי גרף לא יעזור: איזה מאפיין של העומס גורם לoverhead להיות זניחה מלכתחילה.

רמז: היעזרו בקוד ה-torch.cuda.graph מההרצאה. הנקודה שחייבים לתפוס: יצירת tensor חדש בכל צעד מקצה כתובת חדשה ומפרה את הגרף; לכן משתמשים ב-buffer קבוע וכותבים אליו ב-copy_. גרף לא יעזור כשה-kernels גדולים וזמן הdevice שולט - אז הoverhead זניחה ואין מה לרפא. ה-batch הקטן ב-inference הוא בדיוק מה שהופך את ה-kernels לקצרים ואת הoverhead לדומיננטית.