3.2 ניהול זיכרון מהמארח פתרון
פתרון - ניהול זיכרון מהhost¶
הערה: הכתובות, הזמנים המוחלטים ורוחבי הפס בפלטים למטה הם דוגמאות ממכונת H100 SXM עם PCIe Gen5, וישתנו אצלכם לפי הכרטיס, גרסת ה-PCIe ומצב המערכת. מה שלא ישתנה הוא המבנה, סדרי הגודל, ויחסי ה-speedup בין הגרסאות - עליהם התרגיל בנוי. כל הקטעים מתקמפלים עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4 וכדומה).
פתרון תרגיל 1 - pinned מול pageable: מדידת ה-speedup¶
#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)
static float timeH2D(void* dst, const void* src, size_t bytes, int reps) {
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s));
CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaMemcpy(dst, src, bytes, cudaMemcpyHostToDevice)); // warm-up
float best = 1e30f;
for (int r = 0; r < reps; r++) {
CUDA_CHECK(cudaEventRecord(s));
CUDA_CHECK(cudaMemcpy(dst, src, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms; // the best time = the least noise
}
CUDA_CHECK(cudaEventDestroy(s));
CUDA_CHECK(cudaEventDestroy(e));
return best;
}
int main(void) {
int n = 1 << 26; // 67,108,864 elements
size_t bytes = (size_t)n * sizeof(float); // 256 MiB
float* h_pageable = (float*)malloc(bytes);
float* h_pinned;
CUDA_CHECK(cudaMallocHost(&h_pinned, bytes));
for (size_t i = 0; i < (size_t)n; i++) { h_pageable[i] = 1.0f; h_pinned[i] = 1.0f; }
float* d;
CUDA_CHECK(cudaMalloc(&d, bytes));
float tp = timeH2D(d, h_pageable, bytes, 10);
float tn = timeH2D(d, h_pinned, bytes, 10);
double gb = (double)bytes / 1e9;
printf("pageable: %.3f ms -> %.1f GB/s\n", tp, gb / (tp / 1e3));
printf("pinned : %.3f ms -> %.1f GB/s\n", tn, gb / (tn / 1e3));
printf("speedup (pageable/pinned): %.2fx\n", tp / tn);
CUDA_CHECK(cudaFree(d));
CUDA_CHECK(cudaFreeHost(h_pinned));
free(h_pageable);
return 0;
}
הפלט הצפוי (H100 SXM, PCIe Gen5):
למה זה עבד: העברה מזיכרון pageable מחייבת את הדרייבר להעתיק תחילה אל bounce buffer נעול פנימי ורק ממנו לבצע DMA - שתי העתקות, ולכן כמחצית רוחב הפס. זיכרון pinned קבוע פיזית, כך שה-DMA engine קורא ממנו ישירות ומגיע קרוב לרוחב הפס של ה-PCIe Gen5 (כ-64 GB/s תיאורטי). לקיחת הזמן המינימלי על פני עשר הרצות, אחרי חימום, מסננת רעשי תזמון ומיתוג-תהליכים.
איך להכליל: כל buffer שחוזר ומעורב בהעברות H2D/D2H - קלט/פלט בלולאה חמה, staging של pipeline - כדאי לנעול. מדדו תמיד עם cudaEvent ולא עם שעון ה-CPU, וקחו מינימום/חציון על פני חזרות. זכרו את מחיר ה-over-pinning: אל תנעלו זיכרון שאינו מעורב בהעברות.
פתרון תרגיל 2 - חיבור וקטורים עם managed memory¶
__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];
}
int main(void) {
int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
float *a, *b, *c;
CUDA_CHECK(cudaMallocManaged(&a, bytes)); // single pointer for both worlds
CUDA_CHECK(cudaMallocManaged(&b, bytes));
CUDA_CHECK(cudaMallocManaged(&c, bytes));
for (int i = 0; i < n; i++) { a[i] = 1.0f; b[i] = 2.0f; } // write from the host
int bs = 256, nb = (n + bs - 1) / bs;
vecAdd<<<nb, bs>>>(a, b, c, n); // same pointer to the kernel
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize()); // required before reading from the host
double maxErr = 0.0;
for (int i = 0; i < n; i++) { // read from the host - without cudaMemcpy
double d = (double)c[i] - 3.0; if (d < 0) d = -d;
if (d > maxErr) maxErr = d;
}
printf("managed vecadd: max error = %f, cudaMemcpy calls = 0\n", maxErr);
CUDA_CHECK(cudaFree(a)); // cudaFree, not cudaFreeHost
CUDA_CHECK(cudaFree(b));
CUDA_CHECK(cudaFree(c));
return 0;
}
הפלט הצפוי:
למה זה עבד: cudaMallocManaged מחזיר מצביע יחיד שתקף גם על הhost וגם על הdevice. כשהhost כותב ל-a[i], page fault ממקם את העמוד ב-RAM; כשה-kernel ניגש אליו, page fault של ה-GPU מהגר אותו ל-HBM3; וכשהhost קורא את c[i], העמוד מהגר בחזרה. כל ההעברות קורות אוטומטית, ולכן אין ולו cudaMemcpy אחד. ה-cudaDeviceSynchronize הכרחי כי הקריאה מהhost חייבת לחכות שה-kernel יסיים לכתוב.
איך להכליל: managed memory מקצר דרמטית קוד ומצוין לפורטינג של קוד CPU קיים, למבנים עם מצביעים מקוננים, ולפרוטוטייפ מהיר. במחיר: ה-page faults מפוזרים על זמן הריצה, ולכן בקוד ביצועי צריך prefetch (תרגיל 3). שימו לב שמשחררים ב-cudaFree - לא cudaFreeHost, שמיועד רק ל-cudaMallocHost.
פתרון תרגיל 3 - האצת managed memory עם prefetch¶
static float runManaged(int n, bool prefetch) {
size_t bytes = (size_t)n * sizeof(float);
float *a, *b, *c;
CUDA_CHECK(cudaMallocManaged(&a, bytes));
CUDA_CHECK(cudaMallocManaged(&b, bytes));
CUDA_CHECK(cudaMallocManaged(&c, bytes));
int device = 0; CUDA_CHECK(cudaGetDevice(&device));
for (int i = 0; i < n; i++) { a[i] = 1.0f; b[i] = 2.0f; } // data in RAM
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
int bs = 256, nb = (n + bs - 1) / bs;
CUDA_CHECK(cudaEventRecord(s));
if (prefetch) {
CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, device)); // push to the GPU ahead of time
CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, device));
}
vecAdd<<<nb, bs>>>(a, b, c, n);
if (prefetch)
CUDA_CHECK(cudaMemPrefetchAsync(c, bytes, cudaCpuDeviceId)); // back to the host
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
volatile float sink = c[n - 1]; // forces the result to arrive at the host
(void)sink;
float ms = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
CUDA_CHECK(cudaFree(a)); CUDA_CHECK(cudaFree(b)); CUDA_CHECK(cudaFree(c));
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return ms;
}
int main(void) {
for (int n : {1 << 25, 1 << 27}) {
float noPf = runManaged(n, false);
float pf = runManaged(n, true);
printf("n=2^%d no-prefetch: %.3f ms prefetch: %.3f ms speedup: %.2fx\n",
(int)(__builtin_ctz(n)), noPf, pf, noPf / pf);
}
return 0;
}
הפלט הצפוי:
n=2^25 no-prefetch: 6.842 ms prefetch: 2.417 ms speedup: 2.83x
n=2^27 no-prefetch: 27.51 ms prefetch: 9.63 ms speedup: 2.86x
למה זה עבד: בלי prefetch, ה-kernel נאלץ להגר את כל העמודים של a ו-b מ-RAM ל-HBM3 תוך כדי ריצה, עמוד-אחר-עמוד בכל fault - הרבה העברות קטנות ולא יעילות, וגם קריאת c בסוף מהגרת חזרה. עם prefetch, אנחנו ממירים את זה לשלוש העברות גדולות ורציפות שרצות על רוחב הפס המלא של ה-PCIe, בדיוק כמו cudaMemcpy יעיל. ה-speedup גדל עם n כי ככל שיש יותר עמודים, כך overhead של ה-faults הבודדים מצטברת.
איך להכליל: השתמשו ב-managed לנוחות, אבל הוסיפו cudaMemPrefetchAsync בכל נקודה שבה אתם יודעים לאן הנתונים הולכים - לפני kernel שקורא אותם, ואחריו כדי להחזיר תוצאה לhost. לדפוסי גישה מורכבים יותר הוסיפו cudaMemAdvise (למשל SetReadMostly לנתונים שרק נקראים). הכלל: managed מגיע לביצועי ניהול ידני רק כשהעברות ה-migration מקובצות ב-prefetch.
פתרון תרגיל 4 - חפיפת העתקה וחישוב עם שני streams¶
static float serial(float* d_a, float* d_b, float* d_c,
float* h_a, float* h_b, float* h_c, int n) {
size_t bytes = (size_t)n * sizeof(float);
cudaEvent_t s, e; CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
CUDA_CHECK(cudaEventRecord(s));
CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
vecAdd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaEventRecord(e)); CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return ms;
}
static float overlapped(float* d_a, float* d_b, float* d_c,
float* h_a, float* h_b, float* h_c, int n) {
const int nStreams = 2, nChunks = 8;
cudaStream_t st[nStreams];
for (int i = 0; i < nStreams; i++) CUDA_CHECK(cudaStreamCreate(&st[i]));
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 x = st[c % nStreams];
CUDA_CHECK(cudaMemcpyAsync(d_a + off, h_a + off, cb, cudaMemcpyHostToDevice, x));
CUDA_CHECK(cudaMemcpyAsync(d_b + off, h_b + off, cb, cudaMemcpyHostToDevice, x));
vecAdd<<<(chunk + 255) / 256, 256, 0, x>>>(d_a + off, d_b + off, d_c + off, chunk);
CUDA_CHECK(cudaMemcpyAsync(h_c + off, d_c + off, cb, cudaMemcpyDeviceToHost, x));
}
for (int i = 0; i < nStreams; i++) CUDA_CHECK(cudaStreamSynchronize(st[i]));
CUDA_CHECK(cudaEventRecord(e)); CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f; 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 << 25;
size_t bytes = (size_t)n * sizeof(float);
float *h_a, *h_b, *h_c;
CUDA_CHECK(cudaMallocHost(&h_a, bytes)); // pinned - required for overlap
CUDA_CHECK(cudaMallocHost(&h_b, bytes));
CUDA_CHECK(cudaMallocHost(&h_c, bytes));
for (int i = 0; i < n; i++) { h_a[i] = 1.0f; h_b[i] = 2.0f; }
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));
serial(d_a, d_b, d_c, h_a, h_b, h_c, n); // warm-up
float ts = serial(d_a, d_b, d_c, h_a, h_b, h_c, n);
float to = overlapped(d_a, d_b, d_c, h_a, h_b, h_c, n);
double maxErr = 0.0;
for (int i = 0; i < n; i++) { double d = (double)h_c[i] - 3.0; if (d<0) d=-d; if (d>maxErr) maxErr=d; }
printf("serial: %.3f ms overlapped: %.3f ms speedup: %.2fx maxErr=%f\n",
ts, to, ts / to, maxErr);
CUDA_CHECK(cudaFree(d_a)); CUDA_CHECK(cudaFree(d_b)); CUDA_CHECK(cudaFree(d_c));
CUDA_CHECK(cudaFreeHost(h_a)); CUDA_CHECK(cudaFreeHost(h_b)); CUDA_CHECK(cudaFreeHost(h_c));
return 0;
}
הפלט הצפוי:
למה זה עבד: פירוק ל-8 נתחים על שני streams לא-default נותן ל-GPU להריץ במקביל פעולות מ-streams שונים. בזמן שה-DMA engine מעתיק H2D של נתח אחד, מנוע אחר יכול להעתיק D2H של נתח קודם, וה-CUDA Cores מריצים את ה-vecAdd של נתח שלישי. ה-cudaMemcpyAsync באמת אסינכרוני רק מפני שזיכרון הhost pinned - עם pageable הוא היה מסתנכרן ואיבדנו את החפיפה. חיבור וקטורים memory-bound וזול-חישוב, ולכן החפיפה כאן היא בעיקר בין ההעברות בכיוונים ההפוכים, ומכאן ה-speedup המתון של כ-1.5x.
איך להכליל: תבנית ה-chunk-over-streams היא הדרך הסטנדרטית להחביא latency של העברות. ה-speedup תלוי ביחס בין זמן ההעתקה לזמן החישוב לנתח - כשהם דומים, החפיפה מקסימלית (ראו תרגיל בונוס עם kernel כבד). שלושת התנאים שאסור לשכוח: pinned, cudaMemcpyAsync, ו-streams לא-default.
פתרון תרגיל 5 - תזמון kernel מדויק עם cudaEvent¶
int main(void) {
int n = 1 << 27; // 134,217,728
size_t bytes = (size_t)n * sizeof(float);
float *h_a, *h_b;
CUDA_CHECK(cudaMallocHost(&h_a, bytes));
CUDA_CHECK(cudaMallocHost(&h_b, bytes));
for (int i = 0; i < n; i++) { h_a[i] = 1.0f; h_b[i] = 2.0f; }
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));
cudaEvent_t s, e; CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
int bs = 256, nb = (n + bs - 1) / bs;
// H2D time for the two inputs
CUDA_CHECK(cudaEventRecord(s));
CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaEventRecord(e)); CUDA_CHECK(cudaEventSynchronize(e));
float h2d = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&h2d, s, e));
vecAdd<<<nb, bs>>>(d_a, d_b, d_c, n); // warm-up
CUDA_CHECK(cudaDeviceSynchronize());
float best = 1e30f;
for (int r = 0; r < 20; r++) {
CUDA_CHECK(cudaEventRecord(s));
vecAdd<<<nb, bs>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaEventRecord(e)); CUDA_CHECK(cudaEventSynchronize(e));
float ms = 0.0f; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
double bw = (12.0 * n) / (best * 1e6); // GB/s
printf("kernel: %.3f ms -> %.1f GB/s (%.1f%% of 3350)\n",
best, bw, 100.0 * bw / 3350.0);
printf("H2D (2 inputs): %.3f ms -> %.1f GB/s\n",
h2d, (2.0 * bytes) / (h2d * 1e6));
CUDA_CHECK(cudaFree(d_a)); CUDA_CHECK(cudaFree(d_b)); CUDA_CHECK(cudaFree(d_c));
CUDA_CHECK(cudaFreeHost(h_a)); CUDA_CHECK(cudaFreeHost(h_b));
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return 0;
}
הפלט הצפוי:
למה זה עבד: cudaEvent מודד את הזמן על ה-GPU עצמו ברזולוציה תת-מיקרו-שנייה, כך שאנחנו רואים את זמן ה-kernel נטו. חיבור וקטורים memory-bound, ולכן הוא רץ קרוב לרוחב הפס המלא של ה-HBM3 (כ-91% מ-3350 GB/s - ראיה ל-coalescing מצוין). לעומת זאת, זמן ה-H2D של הקלטים גדול פי כ-40 מזמן ה-kernel, כי הוא כבול לרוחב הפס האיטי של ה-PCIe. איטרציית החימום מסירה את overhead ה-JIT וההקצאות מהמדידה.
איך להכליל: תמיד מדדו kernels עם cudaEvent ולא עם שעון CPU, תמיד עם חימום, ותמיד קחו מינימום/חציון על פני חזרות. חשבו רוחב פס אפקטיבי כדי לדעת אם אתם memory-bound (קרוב לתקרת ה-HBM3) או משאירים ביצועים על השולחן. העובדה שה-H2D שולט בזמן הכולל היא בדיוק ההצדקה לכל השיעור: מזערו העברות, נעלו זיכרון, וחפפו את מה שנשאר.
פתרון תרגיל 6 (בונוס) - oversubscription וקרנל כבד¶
// part A - oversubscription: allocation larger than physical memory
__global__ void touch(float* p, size_t n) {
size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
size_t stride = (size_t)blockDim.x * gridDim.x;
for (; i < n; i += stride) p[i] = p[i] + 1.0f;
}
// part B - heavy kernel: lots of compute per element -> better overlap
__global__ void heavy(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float v = a[i] + b[i];
for (int k = 0; k < 256; k++) v = v * 1.0001f + 0.5f; // arithmetic load
c[i] = v;
}
}
הרצת חלק א (על כרטיס עם פחות זיכרון מההקצאה) והפלט הצפוי:
free=23.1 GiB total=23.6 GiB
allocated 30.0 GiB managed (oversubscribed by 6.9 GiB) - OK
touch kernel over-subscribed: 412.7 ms (slow - pages swapping in/out)
הרצת חלק ב (חפיפה עם kernel כבד, מבנה כמו תרגיל 4) והפלט הצפוי:
למה זה עבד: בחלק א, cudaMallocManaged מצליח להקצות 30 GiB על כרטיס עם 24 GiB כי הזיכרון המנוהל הוא וירטואלי - ה-runtime מחליף עמודים בין ה-HBM3 ל-RAM הhost לפי דרישה, בדיוק כמו swap. cudaMalloc היה נכשל כאן ב-out of memory. המחיר: ה-kernel איטי מאוד כי הוא מחכה כל הזמן להחלפות עמודים. בחלק ב, ה-kernel הכבד מבצע 256 פעולות לכל איבר, כך שזמן החישוב לנתח מתקרב לזמן ההעתקה שלו; לפי מודל ה-max(copy, compute) מההרצאה, זה בדיוק המצב שבו החפיפה מקסימלית, ומכאן speedup כמעט פי 2 לעומת ה-1.5x של חיבור הווקטורים הקל.
איך להכליל: oversubscription שימושי כשמבנה הנתונים גדול מה-HBM3 ורק חלק ממנו חם בכל רגע - ה-runtime שומר את החם על ה-GPU. אבל אם הגישה אקראית וכל הנתונים חמים, ההחלפה תחנוק אתכם, ועדיף לפרק את הבעיה ידנית ל-chunks עם streams. לגבי חפיפה: היחס בין copy ל-compute הוא שקובע את הרווח - השקיעו בחפיפה כשהחישוב לנתח דומה בגודלו להעברה שלו, ואל תטרחו כשאחד מהם שולט לגמרי (אז הרכיב האיטי ממילא קובע את הזמן).