7.4 Overhead, latency hiding וחוק Little פתרון
פתרון - Overhead, latency hiding וחוק Little¶
הערה: הזמנים המוחלטים, הoverheads וה-speedups בפלטים למטה הם דוגמאות ממכונת H100 SXM עם CUDA 12, וישתנו אצלכם לפי הכרטיס, גרסת הדרייבר ומצב המערכת. מה שלא ישתנה הוא המבנה, סדרי הגודל והמגמות. תרגילי החשבון (2, 4, 5) הם מדויקים ואינם תלויים בחומרה. כל קטעי הקוד מתקמפלים עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4 וכדומה).
פתרון תרגיל 1 - מדידת overhead הlaunch של 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 noop() { }
int main(void) {
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
const int N = 100000;
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s));
CUDA_CHECK(cudaEventCreate(&e));
// warmup - absorbs context init and one-time JIT
for (int i = 0; i < N; i++) noop<<<1, 1, 0, stream>>>();
CUDA_CHECK(cudaStreamSynchronize(stream));
// measurement - non-default stream
CUDA_CHECK(cudaEventRecord(s, stream));
for (int i = 0; i < N; i++) noop<<<1, 1, 0, stream>>>();
CUDA_CHECK(cudaEventRecord(e, stream));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
printf("stream: %d launches -> %.2f ms (%.2f us/launch)\n",
N, ms, (ms * 1e3) / N);
// measurement - default stream (without the stream parameter)
CUDA_CHECK(cudaEventRecord(s));
for (int i = 0; i < N; i++) noop<<<1, 1>>>();
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
printf("default: %d launches -> %.2f ms (%.2f us/launch)\n",
N, ms, (ms * 1e3) / N);
CUDA_CHECK(cudaEventDestroy(s));
CUDA_CHECK(cudaEventDestroy(e));
CUDA_CHECK(cudaStreamDestroy(stream));
return 0;
}
הפלט הצפוי (H100 SXM):
stream: 100000 launches -> 487.32 ms (4.87 us/launch)
default: 100000 launches -> 501.90 ms (5.02 us/launch)
למה זה עבד: ה-kernel noop ריק לגמרי - אין ארגומנטים, אין חישוב, זמן הdevice שלו כמעט אפס. לכן כמעט כל הזמן שנמדד הוא הhost overhead: הרצף שמאחורי כל <<<...>>> אורז תצורת launch ופונה לדרייבר. חלוקת הזמן הכולל ב-100000 מבודדת את הoverhead לlaunch בודד - כאן ~4.9 מיקרו-שניות, בדיוק בטווח ה-3-10 us שה-glossary מציין. ה-default stream כמעט זהה (מעט איטי יותר בגלל הסמנטיקה המסנכרנת שלו); הפער זעום כי בשני המקרים העלות נשלטת על-ידי אותה עבודת CPU של הנפקת הlaunch.
איך להכליל: זו השיטה למדוד "כמה עולה לי לlaunch" על כל GPU - לולאה של הרבה launches ריקים, חלקי מספרם. את המספר הזה משווים לזמן ההרצה של ה-kernels האמיתיים שלכם: אם הם רצים בזמן דומה לoverhead (יחידות מיקרו-שניות), אתם overhead-bound, וצריך גרפים/מיזוג/קיבוץ. תמיד מדדו עם cudaEvent ולא עם שעון CPU, ותמיד עם חימום.
פתרון תרגיל 2 - חוק Little בעיפרון¶
הנוסחה: concurrency = latency x throughput.
סעיף 1 - פעולות בטיסה.
סעיף 2 - throughput גבוה יותר.
היחס הוא פי 10, בדיוק כמו היחס ב-throughput. התובנה: throughput גבוה יותר דורש מקביליות גבוהה יותר באופן פרופורציוני כדי להישאר מוסתר-latency. מכונה מהירה יותר צריכה יותר עבודה "באוויר" רק כדי להזין את עצמה.
סעיף 3 - מ-פעולות ל-warps.
סעיף 4 - שרשרת תלויה.
סעיף 5 - תרגום ל-occupancy.
לא, לא צריך occupancy גבוה. occupancy צנוע של כ-20% כבר מסתיר לגמרי שרשרת load אופיינית.
למה זה עבד: חוק Little הוא כפל פשוט, אבל היחידות קריטיות. ב-1-2 ספרנו "פעולות בטיסה" (יחידות thread), וב-3-4 המרנו ל-warps בחלוקה ברוחב ה-warp (32), כי ה-warp הוא יחידת ההנפקה. שני ה-13 (מ-400 ומ-416) יוצאים כמעט זהים כי ה-latency בשני המקרים דומה, וזה מאשש שהחשבון עקבי. סעיף 5 הוא ההקדמה לפרק 8: מספר ה-warps הדרוש להסתרה הוא צנוע, ולכן "מקסימום occupancy" אינו המטרה - "מספיק occupancy" הוא.
איך להכליל: לכל latency שאתם רוצים להסתיר, warps ≈ latency_cycles / 32, ו-occupancy ≈ warps / max_warps_per_SM. זה נותן את רצפת ה-occupancy הדרושה. אם ה-latency ארוך יותר (למשל שרשרת של כמה loads תלויים) - צריך יותר warps; אם חושפים ILP - צריך פחות (תרגיל 4).
פתרון תרגיל 3 - המרת הלולאה לגרף וכימות הoverhead שהוסרה¶
static float timeLoop(cudaStream_t stream, int N, int reps) {
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
for (int i = 0; i < N; i++) noop<<<1,1,0,stream>>>(); // warmup
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++) noop<<<1,1,0,stream>>>();
CUDA_CHECK(cudaEventRecord(e, stream));
CUDA_CHECK(cudaEventSynchronize(e));
float ms; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return best;
}
static float timeGraph(cudaStream_t stream, int N, int reps) {
cudaGraph_t graph; cudaGraphExec_t ge;
CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
for (int i = 0; i < N; i++) noop<<<1,1,0,stream>>>();
CUDA_CHECK(cudaStreamEndCapture(stream, &graph));
CUDA_CHECK(cudaGraphInstantiate(&ge, graph, 0));
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaGraphLaunch(ge, stream)); // warmup
CUDA_CHECK(cudaStreamSynchronize(stream));
float best = 1e30f;
for (int r = 0; r < reps; r++) {
CUDA_CHECK(cudaEventRecord(s, stream));
CUDA_CHECK(cudaGraphLaunch(ge, stream));
CUDA_CHECK(cudaEventRecord(e, stream));
CUDA_CHECK(cudaEventSynchronize(e));
float ms; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
CUDA_CHECK(cudaGraphExecDestroy(ge)); CUDA_CHECK(cudaGraphDestroy(graph));
return best;
}
int main(void) {
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
const int N = 512;
float lp = timeLoop(stream, N, 50);
float gr = timeGraph(stream, N, 50);
printf("N=%d loop: %.3f ms graph: %.3f ms speedup: %.2fx\n",
N, lp, gr, lp / gr);
printf("overhead removed per run: %.3f ms\n", lp - gr);
CUDA_CHECK(cudaStreamDestroy(stream));
return 0;
}
הפלט הצפוי:
למה זה עבד: בלולאה, הhost שילם ~4.9 מיקרו-שניות עבור כל אחד מ-512 הlaunches - סה"כ ~2.5 ms. בגרף, כל 512 הlaunches נלכדו כצמתים, ה-instantiate עשה את העבודה היקרה פעם אחת, וכל cudaGraphLaunch הוא קריאת host אחת שמגישה את כולן - סה"כ ~0.18 ms. הגרף לא האיץ שום kernel (הם ריקים ממילא, זמן הdevice שלהם אפס); הוא הסיר רק את הhost overhead, וזה בדיוק ~2.33 ms שהתפנו. זהו הניסוי הנקי ביותר לבידוד overhead הlaunch, כי אין שום חישוב שיערבב.
איך להכליל: overhead_removed ≈ (launches x overhead_per_launch) - replay_cost. הרווח גדל עם מספר ה-kernels הקטנים (ראו את סריקת ה-N בשיעור 5.5). כשה-kernels ריקים או זעירים, הגרף הוא ניצחון ברור; כשהם גדולים וזמן הdevice שולט, אין overhead משמעותית לרפא.
פתרון תרגיל 4 - חישוב ה-occupancy הדרוש לתרחיש נתון¶
התרחיש: שני loads בלתי-תלויים (400 מחזורים כל אחד) + שרשרת אריתמטית תלויה של 40 מחזורים.
סעיף 1 - ILP של warp בודד. מכיוון ששני ה-loads אינם תלויים זה בזה, ה-warp מנפיק את שניהם לפני שהוא נתקע. לכן warp בודד מחזיק 2 פעולות בטיסה (ILP = 2).
סעיף 2 - latency אפקטיבי להסתרה. שני ה-loads חופפים בזמן (הם בטיסה יחד), ולכן ה-latency שנחשף מהם הוא 400 מחזורים (לא 800), ואחריו 40 מחזורי אריתמטיקה שתלויים בהם:
סעיף 3 - warps ו-occupancy. כלל האצבע, מורחב ל-ILP: warps ≈ latency / (32 x ILP):
סעיף 4 - התרחיש התלוי. אם שני ה-loads תלויים זה בזה (pointer-chasing, שרשרת אחת), הם רצים בטור וה-ILP צונח ל-1:
critical path = 400 + 400 + 40 = 840 cycles, ILP = 1
warps = 840 / 32 ≈ 26 warps
occupancy = 26 / 64 ≈ 0.41 ≈ 41%
למה זה עבד: ILP ו-occupancy הם שני מנופים חלופיים לאותו יעד - להחזיק מספיק פעולות בטיסה כדי לספק את חוק Little. במקרה הבלתי-תלוי, כל warp כבר תורם 2 פעולות בטיסה, ולכן צריך פחות warps (7, occupancy 11%). במקרה התלוי, כל warp תורם פעולה אחת בלבד, וגם ה-latency הכולל גדל (הגישות בטור), ולכן צריך פי כמעט 4 יותר warps (26, occupancy 41%). אותו kernel, שני דפוסי גישה - דרישת occupancy שונה לחלוטין.
איך להכליל: כשאתם קצרים ב-warps (occupancy נמוך בגלל register pressure או shared memory), חשפו ILP - סדרו את הקוד כך שיהיו פעולות בלתי-תלויות ברצף (unrolling, loads מוקדמים) - וזה מפחית את מספר ה-warps הדרוש להסתרה. הנוסחה המאחדת: warps ≈ latency / (32 x ILP).
פתרון תרגיל 5 - הסבר תוצאת Volkov (30 מול 24)¶
סעיף 1 - למה זה מפתיע. ה-latency של גישה לזיכרון גלובלי הוא מאות מחזורים; ה-latency של פעולה אריתמטית הוא יחידות מחזורים - פער של פי כמה עשרות. הציפייה הנאיבית: כדי להסתיר latency גדול פי כמה עשרות, צריך פי כמה עשרות יותר warps. בפועל המספרים כמעט זהים (30 מול 24).
סעיף 2 - חוק Little לשני המקרים.
memory: concurrency = high latency x low throughput
arithmetic: concurrency = low latency x high throughput
סעיף 3 - למה המכפלה דומה. ה-concurrency הוא המכפלה של latency ו-throughput, לא ה-latency לבדו. לזיכרון latency גבוה אבל throughput נמוך (רוחב-פס ה-HBM נמוך יחסית); לאריתמטיקה latency נמוך אבל throughput גבוה (רוחב-פס אריתמטי גבוה). כל אחד מפצה על הגורם הגדול שלו בגורם קטן, והמכפלות מתקזזות לכדי concurrency דומה - ולכן מספר warps דומה.
סעיף 4 - הקשר ל"קיר הזיכרון". הגורם שthroughput בחוק Little הוא שמשקף את הפער: רוחב-הפס האריתמטי גבוה בהרבה מרוחב-פס הזיכרון (היחס Tensor Core:CUDA Core הוא כ-100:1, וקיר-הזיכרון מ-7.1 אומר בדיוק שהחישוב מתקדם מהר יותר מהזיכרון). ה-throughput הנמוך של הזיכרון הוא מה שממתן את ה-concurrency הדרוש למרות ה-latency הגבוה.
סעיף 5 - למה זו תכונה שימושית. מכיוון שאותו מספר warps (~24-30) מסתיר גם latency אריתמטי וגם latency של זיכרון, אותה תצורת occupancy מסתירה עומס מעורב של חישוב וזיכרון - וזה בדיוק מה שרוב ה-kernels עושים. המכונה מתוכננת כך שמאזן ה-warps עובד לשני סוגי ה-latency בבת אחת.
למה זה עבד: המלכודת האינטואיטיבית היא להסתכל על latency לבד. חוק Little מכריח להכפיל ב-throughput, ושם הקיזוז מתגלה. זו דוגמה מושלמת לכך שהמדד הנכון להסתרה הוא concurrency (המכפלה), לא latency.
איך להכליל: בכל פעם שאתם מנתחים כמה מקביליות צריך, אל תסתכלו על latency לבד - הכפילו ב-throughput של אותו משאב. משאב איטי (זיכרון) עם throughput נמוך לא בהכרח דורש יותר מקביליות ממשאב מהיר (אריתמטיקה) עם throughput גבוה.
פתרון תרגיל 6 (בונוס) - לראות את הסתרת ה-latency קורית¶
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) \
do { cudaError_t e_=(call); if(e_!=cudaSuccess){ \
fprintf(stderr,"CUDA %s:%d %s\n",__FILE__,__LINE__, \
cudaGetErrorString(e_)); exit(1);} } while(0)
// every thread chases a dependent pointer chain - memory latency dominates
__global__ void chase(const int* next, int* out, int steps) {
int p = threadIdx.x;
for (int s = 0; s < steps; s++) p = next[p];
out[threadIdx.x] = p;
}
int main(void) {
const int MAXT = 1024; // up to 32 warps per block
const size_t NELEM = 64ull << 20; // 64M ints = 256MB, much larger than L2 (50MB) => every load misses to HBM
const size_t STRIDE = 1048583; // odd => a single cycle across the whole array
const int STEPS = 500000; // enough steps for the traversal to exceed L2 capacity and truly miss
int* h_next = (int*)malloc(NELEM * sizeof(int));
for (size_t i = 0; i < NELEM; i++) // a pointer chain that jumps ~4MB per step - every load misses to global memory
h_next[i] = (int)((i + STRIDE) % NELEM);
int *d_next, *d_out;
CUDA_CHECK(cudaMalloc(&d_next, NELEM * sizeof(int)));
CUDA_CHECK(cudaMalloc(&d_out, MAXT * sizeof(int)));
CUDA_CHECK(cudaMemcpy(d_next, h_next, NELEM * sizeof(int), cudaMemcpyHostToDevice));
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
printf("%6s %8s %12s\n", "warps", "ms", "Gacc/s");
for (int W = 1; W <= 32; W *= 2) {
int threads = W * 32;
chase<<<1, threads>>>(d_next, d_out, STEPS); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
float best = 1e30f;
for (int r = 0; r < 20; r++) {
CUDA_CHECK(cudaEventRecord(s));
chase<<<1, threads>>>(d_next, d_out, STEPS);
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
float ms; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
double acc = (double)threads * STEPS; // total accesses
printf("%6d %8.3f %12.3f\n", W, best, acc / (best * 1e6));
}
CUDA_CHECK(cudaFree(d_next)); CUDA_CHECK(cudaFree(d_out));
free(h_next);
return 0;
}
הפלט הצפוי (H100 SXM, block בודד):
warps ms Gacc/s
1 118.700 0.135
2 119.200 0.268
4 120.100 0.533
8 122.800 1.042
16 130.500 1.962
32 170.000 3.012
למה זה עבד: כל thread רודף שרשרת מצביעים תלויה על פני מערך של 256MB - גדול בהרבה מנפח ה-L2 - ולכן כל load מחטיא באמת לזיכרון הגלובלי ונתקע ~400 מחזורים. עם warp אחד בלבד, ה-warp scheduler נשאר ללא warp כשיר בזמן ההמתנה - הזמן כמעט קבוע אבל מעט מאוד גישות הושלמו (0.14 Gacc/s). ככל שמוסיפים warps, יש יותר עבודה בלתי-תלויה להנפיק בזמן שאחרים תקועים: הזמן כמעט לא גדל (מ-118.7 ל-130.5 ms מ-1 ל-16 warps), אבל מספר הגישות גדל פי 16 - ולכן ה-throughput מזנק כמעט לינארית. סביב 16-32 warps ה-latency מוסתר במלואו וה-throughput מתחיל להתרווח (הזמן מתחיל לגדול משמעותית - מ-130.5 ל-170.0 ms - כי כבר אין latency להסתיר, ומעכשיו מוסיפים עבודה אמיתית). נקודת ההתרווחות סביב 16-24 warps תואמת יפה את תוצאת Volkov (~30 warps לזיכרון) ואת החישוב התיאורטי מתרגיל 2.
איך להכליל: זו ההדגמה החיה של חוק Little ושל "מספיק occupancy, לא מקסימלי". מוסיפים warps כל עוד ה-throughput עולה; ברגע שהעקומה מתרווחת - הגעתם למספר ה-warps שמסתיר את ה-latency, ומעבר לזה occupancy נוסף לא קונה throughput. אם תריצו ncu --set full ./overhead, תראו את אותו סיפור במדדים: "Achieved Occupancy" עולה, וסיבת ה-stall הדומיננטית Stall Long Scoreboard (המתנה ל-load מהזיכרון הגלובלי) נעשית מכוסה יותר ויותר ככל שיש warps שמחפים עליה.