5.5 CUDA Graphs לצמצום overhead פתרון
פתרון - CUDA Graphs לצמצום overhead¶
הערה: הזמנים המוחלטים, הoverheads לכל launch וה-speedups בפלטים למטה הם דוגמאות ממכונת H100 SXM עם CUDA 12 ו-PCIe Gen5, וישתנו אצלכם לפי הכרטיס, גרסת הדרייבר ומצב המערכת. מה שלא ישתנה הוא המבנה, סדרי הגודל, והמגמות - שזמן הלולאה לינארי במספר הlaunches בעוד שזמן ה-replay כמעט קבוע, ולכן ה-speedup גדל עם מספר ה-kernels. כל הקטעים מתקמפלים עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4 וכדומה).
פתרון תרגיל 1 - מדידת overhead הlaunch בלולאה¶
#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; }
// runs N launches in a loop, measures with cudaEvent, returns milliseconds
static float timeLoop(float* d, cudaStream_t stream, int N, int reps) {
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s));
CUDA_CHECK(cudaEventCreate(&e));
// warm-up
for (int i = 0; i < N; i++) tiny<<<1, 1, 0, stream>>>(d);
CUDA_CHECK(cudaStreamSynchronize(stream));
float best = 1e30f;
for (int r = 0; r < reps; r++) {
CUDA_CHECK(cudaEventRecord(s, stream));
for (int i = 0; i < N; i++) tiny<<<1, 1, 0, stream>>>(d);
CUDA_CHECK(cudaEventRecord(e, stream));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
CUDA_CHECK(cudaEventDestroy(s));
CUDA_CHECK(cudaEventDestroy(e));
return best;
}
int main(void) {
float* d;
CUDA_CHECK(cudaMalloc(&d, sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
const int N = 512;
float loop_ms = timeLoop(d, stream, N, 50);
printf("loop: %d launches -> %.3f ms (%.2f us/launch)\n",
N, loop_ms, (loop_ms * 1e3) / N);
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d));
return 0;
}
הפלט הצפוי (H100 SXM):
למה זה עבד: ה-kernel tiny עושה עבודה זניחה (חיבור אחד), כך שכמעט כל הזמן שנמדד הוא overhead הhost - אריזת הארגומנטים, תצורת הlaunch והפנייה לדרייבר עבור כל אחת מ-512 הlaunches. חלוקת הזמן הכולל ב-512 מבודדת את הoverhead לlaunch בודדת (כאן ~5.56 מיקרו-שניות), בדיוק בטווח ה-3-10 us שהזכרנו בהרצאה. לקיחת המינימום על פני 50 חזרות, אחרי חימום, מסננת רעשי תזמון ומיתוג-תהליכים ואת ה-JIT החד-פעמי.
איך להכליל: זו השיטה הכללית למדוד "כמה עולה לי להשיק" - לולאה של הרבה kernels זעירים, חלקי מספרם. אם המספר שאתם מקבלים דומה בסדר גודל לזמן ההרצה של ה-kernels האמיתיים שלכם, אתם overhead-bound, וגרף CUDA הוא התיקון. תמיד מדדו עם cudaEvent (זמן GPU) ולא עם שעון CPU, תמיד עם חימום, ותמיד קחו מינימום/חציון.
פתרון תרגיל 2 - לכידת אותו רצף לגרף ומדידת הריפוי¶
// captures N tiny launches into a graph, measures a single replay, returns milliseconds
static float timeGraph(float* d, cudaStream_t stream, int N, int reps) {
cudaGraph_t graph;
cudaGraphExec_t graphExec;
CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
for (int i = 0; i < N; i++) tiny<<<1, 1, 0, stream>>>(d);
CUDA_CHECK(cudaStreamEndCapture(stream, &graph));
CUDA_CHECK(cudaGraphInstantiate(&graphExec, graph, 0));
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaGraphLaunch(graphExec, stream)); // warm-up
CUDA_CHECK(cudaStreamSynchronize(stream));
float best = 1e30f;
for (int r = 0; r < reps; r++) {
CUDA_CHECK(cudaEventRecord(s, stream));
CUDA_CHECK(cudaGraphLaunch(graphExec, stream));
CUDA_CHECK(cudaEventRecord(e, stream));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
CUDA_CHECK(cudaGraphExecDestroy(graphExec));
CUDA_CHECK(cudaGraphDestroy(graph));
return best;
}
int main(void) {
float* d;
CUDA_CHECK(cudaMalloc(&d, sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
const int N = 512;
float loop_ms = timeLoop(d, stream, N, 50);
float graph_ms = timeGraph(d, stream, N, 50);
printf("loop: %.3f ms graph: %.3f ms speedup: %.2fx\n",
loop_ms, graph_ms, loop_ms / graph_ms);
// correctness check: reset, run a known replay, verify
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaGraph_t g; cudaGraphExec_t ge;
CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
for (int i = 0; i < N; i++) tiny<<<1,1,0,stream>>>(d);
CUDA_CHECK(cudaStreamEndCapture(stream, &g));
CUDA_CHECK(cudaGraphInstantiate(&ge, g, 0));
const int R = 100;
for (int r = 0; r < R; r++) CUDA_CHECK(cudaGraphLaunch(ge, stream));
CUDA_CHECK(cudaStreamSynchronize(stream));
float host = 0.0f;
CUDA_CHECK(cudaMemcpy(&host, d, sizeof(float), cudaMemcpyDeviceToHost));
printf("verify: d[0]=%.0f (expected %d)\n", host, N * R);
CUDA_CHECK(cudaGraphExecDestroy(ge)); CUDA_CHECK(cudaGraphDestroy(g));
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d));
return 0;
}
הפלט הצפוי:
למה זה עבד: בלכידה, 512 הlaunches לא רצו - הן הוקלטו כצמתים בגרף. ה-cudaGraphInstantiate עשה פעם אחת את כל העבודה היקרה (אימות ה-DAG, מיפוי משאבים, הכנת מבני הפקודות), וכל cudaGraphLaunch אחריו הוא קריאת host אחת שמגישה את כל ה-512 בבת אחת. לכן זמן ה-replay צנח מ-2.85 ms (512 launches נפרדות) ל-0.32 ms - כמעט פי 9. שימו לב שהגרף לא האיץ את ה-kernels עצמם; הוא רק הסיר את overhead הhost, וזה כל הסיפור כשה-kernels זעירים. האימות מראה נכונות מלאה: 512 launches כפול 100 replays נותנות 51200.
איך להכליל: זהו הדפוס הסטנדרטי - לכידה פעם אחת, instantiate פעם אחת, ואז replay חוזר בלולאה החמה. שמרו את ה-graphExec בין הרצות ואל תלכדו מחדש בכל איטרציה (הלכידה וה-instantiation יקרים; רק ה-replay זול). זכרו שבתוך הלכידה אסורות פעולות חוסמות, ושחייבים stream לא-default.
פתרון תרגיל 3 - סריקת מספר הlaunches¶
int main(void) {
float* d;
CUDA_CHECK(cudaMalloc(&d, sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
int Ns[] = {8, 64, 512, 4096};
printf("%6s %10s %10s %8s\n", "N", "loop(ms)", "graph(ms)", "speedup");
for (int N : Ns) {
float lp = timeLoop(d, stream, N, 50);
float gr = timeGraph(d, stream, N, 50);
printf("%6d %10.3f %10.3f %7.2fx\n", N, lp, gr, lp / gr);
}
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d));
return 0;
}
הפלט הצפוי:
N loop(ms) graph(ms) speedup
8 0.046 0.021 2.19x
64 0.357 0.048 7.44x
512 2.847 0.318 8.95x
4096 22.71 2.41 9.42x
למה זה עבד: זמן הלולאה גדל לינארית ב-N - כל launch נוספת מוסיפה את הoverhead הקבועה שלה (~5.5 us), ולכן 4096 launches עולות בערך פי 8 מ-512. זמן ה-replay גם הוא גדל, אבל לא בגלל overhead host (שקבועה - קריאה אחת) אלא בגלל שזמן ההרצה של יותר kernels על הdevice גדל. מכיוון שהמונה (loop) גדל מהר יותר מהמכנה (graph), ה-speedup עולה עם N. עבור N=8 הגרף כמעט לא עוזר - שמונה launches הן overhead זניחה, ואין הרבה מה לרפא.
איך להכליל: הכלל שנחלץ מהטבלה הוא ליבת השיעור - הרווח מגרף גדל ככל שיש יותר kernels קטנים. אם העומס שלכם משגר מאות kernels קצרים לאיטרציה (כמו LLM decode), הגרף ייתן שיפור גדול; אם הוא משגר מעט kernels גדולים, אל תטרחו - הoverhead שם ממילא זניחה. השתמשו במדידה כזו כדי להחליט אם משתלם להשקיע בגרף.
פתרון תרגיל 4 - המגבלה: מצביעים צרובים¶
int main(void) {
float* d;
CUDA_CHECK(cudaMalloc(&d, sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
// capture a graph that writes to d
cudaGraph_t g; cudaGraphExec_t ge;
CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
for (int i = 0; i < 512; i++) tiny<<<1,1,0,stream>>>(d);
CUDA_CHECK(cudaStreamEndCapture(stream, &g));
CUDA_CHECK(cudaGraphInstantiate(&ge, g, 0));
CUDA_CHECK(cudaGraphLaunch(ge, stream)); // valid replay
CUDA_CHECK(cudaStreamSynchronize(stream));
printf("captured d = %p\n", (void*)d);
// now free d and allocate a new d2
CUDA_CHECK(cudaFree(d));
float* d2;
CUDA_CHECK(cudaMalloc(&d2, sizeof(float)));
CUDA_CHECK(cudaMemset(d2, 0, sizeof(float)));
printf("new d2 = %p\n", (void*)d2);
// run the graph again - it still writes to d's old address!
CUDA_CHECK(cudaGraphLaunch(ge, stream));
CUDA_CHECK(cudaStreamSynchronize(stream));
float host = 0.0f;
CUDA_CHECK(cudaMemcpy(&host, d2, sizeof(float), cudaMemcpyDeviceToHost));
printf("d2[0] after replay = %.1f (d2 reused the freed address -> graph silently overwrote it)\n", host);
CUDA_CHECK(cudaGraphExecDestroy(ge)); CUDA_CHECK(cudaGraphDestroy(g));
CUDA_CHECK(cudaFree(d2));
CUDA_CHECK(cudaStreamDestroy(stream));
return 0;
}
הפלט הצפוי (הרצה רגילה):
captured d = 0x7f2c4a000000
new d2 = 0x7f2c4a000000
d2[0] after replay = 512.0 (d2 reused the freed address -> graph silently overwrote it)
והרצה תחת compute-sanitizer ./graphs (כשה-allocator מחזיר כתובת שונה ל-d2, או כשה-replay נוגע בזיכרון משוחרר):
========= COMPUTE-SANITIZER
captured d = 0x7f2c4a000000
new d2 = 0x7f2c4a000200
========= Invalid __global__ write of size 4 bytes
========= at tiny(float*)+0x70
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x7f2c4a000000 is out of bounds
========= and is 512 bytes before the nearest allocation
=========
========= ERROR SUMMARY: 512 errors
למה זה עבד: הכתובת של d נצרבה לתוך graphExec בזמן הלכידה כארגומנט קבוע. אחרי cudaFree(d), ה-replay ממשיך לכתוב לאותה כתובת ישנה. אם ה-allocator במקרה החזיר את אותה כתובת ל-d2 (כמו בהרצה הראשונה למעלה, כי CUDA ממחזר את הבלוק ששוחרר זה עתה), התוצאה נראית "כאילו עבדה" - אבל זה מזל, לא נכונות: הגרף כתב ל-0x...000000 בלי לדעת דבר על d2. אם ה-allocator החזיר כתובת אחרת, הגרף כותב לזיכרון משוחרר - use-after-free ש-compute-sanitizer חושף כ-Invalid global write. זו הדגמה מדויקת למגבלת "המצביעים הצרובים" מההרצאה.
איך להכליל: לעולם אל תשחררו ותקצו מחדש buffer שגרף חי משתמש בו. אם אתם חייבים להפנות את הגרף ל-buffer חדש, הדרך הנכונה היא cudaGraphExecUpdate(graphExec, newGraph, ...) - לוכדים מחדש גרף (זול, בלי instantiate מלא) עם המצביעים החדשים ומעדכנים את הגרף בר-ההרצה הקיים כל עוד הטופולוגיה זהה. ב-PyTorch המקבילה היא לכתוב לתוך ה-buffer הסטטי הקבוע עם copy_ במקום ליצור tensor חדש. הכלל: יציבות הקצאות היא תנאי נכונות, לא רק ביצועים.
פתרון תרגיל 5 - השוואה ל-cudaEvent על ה-CPU מול ה-GPU¶
#include <chrono>
int main(void) {
float* d;
CUDA_CHECK(cudaMalloc(&d, sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, sizeof(float)));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
const int N = 512;
// warm-up
for (int i = 0; i < N; i++) tiny<<<1,1,0,stream>>>(d);
CUDA_CHECK(cudaStreamSynchronize(stream));
// measure GPU with cudaEvent
float gpu_ms = timeLoop(d, stream, N, 50);
// measure CPU issue only (no sync in between)
auto t0 = std::chrono::steady_clock::now();
for (int i = 0; i < N; i++) tiny<<<1,1,0,stream>>>(d);
auto t1 = std::chrono::steady_clock::now();
double cpu_issue_ms = std::chrono::duration<double, std::milli>(t1 - t0).count();
CUDA_CHECK(cudaStreamSynchronize(stream)); // only after the measurement
// measure CPU including sync
auto t2 = std::chrono::steady_clock::now();
for (int i = 0; i < N; i++) tiny<<<1,1,0,stream>>>(d);
CUDA_CHECK(cudaStreamSynchronize(stream));
auto t3 = std::chrono::steady_clock::now();
double cpu_full_ms = std::chrono::duration<double, std::milli>(t3 - t2).count();
printf("GPU (cudaEvent): %.3f ms\n", gpu_ms);
printf("CPU issue-only (no sync): %.3f ms\n", cpu_issue_ms);
printf("CPU issue+sync: %.3f ms\n", cpu_full_ms);
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d));
return 0;
}
הפלט הצפוי:
למה זה עבד: כאן ה-kernels כה זעירים שהhost הוא צוואר הבקבוק - הוא בקושי מספיק לדחוף launch אחת לתור לפני שה-GPU כבר סיים אותה. לכן שלוש המדידות קרובות: זמן ההנפקה של הhost (CPU issue-only) כמעט שווה לזמן ה-GPU, כי ה-GPU רץ צמוד להנפקה. במקרה כללי עם kernels כבדים יותר, CPU issue-only היה קטן משמעותית מזמן ה-GPU: הhost היה מנפיק את כל 512 הlaunches במהירות וממשיך הלאה, בעוד שה-GPU עדיין טוחן. הפער בין issue-only ל-issue+sync הוא בדיוק כמות העבודה שה-GPU עשה אחרי שהhost סיים להנפיק.
איך להכליל: cudaEvent מודד תמיד את זמן ה-GPU, ולכן הוא הכלי הנכון להשוואת ביצועי kernels וגרפים. שעון ה-CPU בלי סנכרון מודד את "overhead ההנפקה" - כמה עמוס הhost, וזה בדיוק המדד שגרף משפר. כשהשניים קרובים, אתם host-bound (overhead-bound) והגרף יעזור מאוד; כשזמן ה-GPU גדול בהרבה מזמן ההנפקה, אתם device-bound והגרף כמעט לא יזיז.
פתרון תרגיל 6 (בונוס) - סקיצת torch.cuda.graph לצעד inference¶
import torch
model = build_model().cuda().eval()
# 1. static buffers - this is where the addresses the graph will use get baked in.
# do not replace them with a new tensor after capture.
static_input = torch.zeros(1, seq_len, hidden, device="cuda")
static_output = None
# 2. warm-up on a side stream - required before capture (stabilizes allocations, runs JIT)
s = torch.cuda.Stream()
s.wait_stream(torch.cuda.current_stream())
with torch.cuda.stream(s):
for _ in range(3):
static_output = model(static_input)
torch.cuda.current_stream().wait_stream(s)
# 3. capture - the whole forward pass (hundreds of kernels) becomes a single runnable graph
g = torch.cuda.CUDAGraph()
with torch.cuda.graph(g):
static_output = model(static_input)
# 4. repeated run per token/request - one launch instead of hundreds
@torch.inference_mode()
def infer(new_input):
static_input.copy_(new_input) # write into the fixed address - not a new tensor!
g.replay() # one host call for the whole forward pass
return static_output.clone() # copy out before the next step overwrites it
הסבר הסימונים המבוקשים:
Where the addresses get baked in: at step 3 (capture). The addresses of static_input,
of the weights, and of all the model's intermediate buffers - are all baked into the graph.
Why static_input is fixed: the graph writes/reads from its baked-in address. A new tensor
= a new address = the graph ignores it and reads from the old one (exactly exercise 4).
Why copy_ and not a new tensor: copy_ writes into the same buffer without changing the address;
a new new_input every step would violate the graph.
למה זה עבד: צעד decode של LLM מריץ את אותו forward pass של השכבות שוב ושוב לכל טוקן, עם אותם משקלים ואותם buffers - טופולוגיה קבועה לחלוטין, בדיוק מה שגרף דורש. מההרצאה: כ-480 launches לטוקן כפול ~5 מיקרו-שניות ≈ 2.4 מילישניות של host overhead בלבד לכל טוקן. הגרף ממיר את 480 הlaunches לlaunch אחת, ומחזיר את התקציב הזה כמעט במלואו - קריטי כשיעד ה-latency הוא מילישניות בודדות. ה-copy_ וה-buffers הסטטיים הם התרגום הישיר של מגבלת "יציבות ההקצאות" לעולם ה-tensors.
איך להכליל: השתמשו ב-torch.cuda.graph (או ב-torch.cuda.make_graphed_callables למקרים פשוטים) בכל עומד inference עם צורה קבועה וחוזרת. גרף לא יעזור כשה-kernels גדולים וזמן הdevice שולט - למשל training ב-batch ענק, שם כל kernel רץ מילישניות והoverhead זניחה. המאפיין שהופך את הoverhead לדומיננטית הוא kernels קצרים ורבים, וזה בדיוק מה שקורה ב-inference עם batch קטן: פחות עבודה לכל kernel, אותו מספר launches, ולכן הoverhead תופסת נתח גדול - והגרף מרפא אותה.