3.2 ניהול זיכרון מהמארח תרגול
תרגול - ניהול זיכרון מהhost¶
בתרגול הזה תמדדו ותרגישו בעצמכם את כל ארבעת הכלים מההרצאה. תתחילו במדידת ההפרש בין העברה מזיכרון pageable להעברה מזיכרון pinned ותדווחו על ה-speedup, תמירו את חיבור הווקטורים ל-cudaMallocManaged ותראו שהוא "פשוט עובד" ואז תאיצו אותו עם prefetch, תחפפו העתקה עם חישוב בעזרת שני streams וזיכרון pinned ותמדדו כמה חסכתם, ולבסוף תתזמנו kernel בדיוק עם cudaEvent. עבדו לפי הסדר - כל תרגיל מוסיף כלי אחד. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA (רק המספרים המוחלטים ישתנו).
הכנה¶
השתמשו בבסיס משיעור 0.3: אותו מקרו CUDA_CHECK, אותו kernel vecAdd. צרו קובץ memmgmt.cu והתחילו ממנו. הקומפילציה לאורך כל התרגול (החליפו את sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):
מקרו הבדיקה שנשתמש בו לכל אורך הדרך:
#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)
תרגיל 1 - pinned מול pageable: מדידת ה-speedup¶
- הקצו שני buffers של host בגודל
1 << 26איברים (256 MiB): אחד רגיל עםmalloc(pageable), ואחד עםcudaMallocHost(pinned). הקצו buffer יעד אחד על הdevice עםcudaMalloc. - מדדו עם
cudaEventאת זמן ה-H2D מהמקור ה-pageable אל הdevice. חזרו על המדידה 10 פעמים וקחו את הזמן החציוני (או הממוצע ללא ההרצה הראשונה). - מדדו באותה דרך את זמן ה-H2D מהמקור ה-pinned.
- חשבו לכל אחד את רוחב הפס האפקטיבי ב-GB/s:
bytes / (ms * 1e6), ודווחו את ה-speeduptime_pageable / time_pinned. - שחררו נכון:
freeל-pageable,cudaFreeHostל-pinned,cudaFreeלdevice.
רמז: הריצו העברת חימום אחת לפני לולאת המדידה - ה-context, ההקצאות והדרייבר מתאתחלים בפעם הראשונה. השתמשו ב-cudaEventRecord לפני ואחרי ה-cudaMemcpy, ואז cudaEventSynchronize(stop) ו-cudaEventElapsedTime. אתם אמורים לראות speedup של פי 2 עד פי 3 לפחות; ה-pinned אמור להתקרב לרוחב הפס של ה-PCIe (Gen5: 50-55 GB/s; Gen4: 20-26 GB/s).
תרגיל 2 - חיבור וקטורים עם managed memory¶
- שכתבו את חיבור הווקטורים כך שכל שלושת המערכים
a,b,cיוקצו ב-cudaMallocManagedבמקום ב-malloc+cudaMalloc. - אתחלו את
aו-bישירות מהhost על אותם מצביעים (בלי שוםcudaMemcpy), השיקו אתvecAddעל אותם מצביעים, ואחריcudaDeviceSynchronizeקראו אתcישירות מהhost ואמתו שהתוצאה נכונה. - שימו לב כמה שורות קוד נעלמו לעומת גרסת 0.3. ספרו: כמה קריאות
cudaMemcpyיש עכשיו? - שחררו את הכל עם
cudaFree(לאcudaFreeHost).
רמז: אין cudaMemcpy בכלל - זו כל הנקודה. המצביע היחיד תקף בשני העולמות, וה-runtime מהגר את העמודים לבד ב-page faults. ודאו ש-cudaDeviceSynchronize מופיע לפני שאתם קוראים את c מהhost, אחרת תקראו לפני שה-kernel סיים לכתוב.
תרגיל 3 - האצת managed memory עם prefetch¶
- הוסיפו מדידת
cudaEventמסביב לכל האזור של גרסת ה-managed מתרגיל 2: מרגע שה-kernel מושק ועד שהתוצאה זמינה בhost (כלומר כולל את זמן ה-page faults). כדי לבודד את זה נקי, געו בנתונים בhost לפני המדידה כך שיהיו ב-RAM, ואז מדדו את ה-kernel + הגישה מהhost לתוצאה. - עכשיו הוסיפו
cudaMemPrefetchAsync(a, bytes, device)ו-cudaMemPrefetchAsync(b, bytes, device)לפני הlaunch, ו-cudaMemPrefetchAsync(c, bytes, cudaCpuDeviceId)אחרי הlaunch. מדדו שוב. - השוו את שני הזמנים. בכמה השתפר?
- הריצו עם
n = 1 << 25וגםn = 1 << 27וראו איך הפער גדל עם גודל הנתונים.
רמז: int device = 0; cudaGetDevice(&device);. ה-prefetch ממיר "הרבה faults קטנים בזמן ה-kernel" ל"העברה אחת גדולה ורציפה" לפני ה-kernel. אם אינכם רואים הבדל, ודאו שאתם באמת גורמים ל-faults - כלומר שהעמודים נמצאים בצד הלא-נכון לפני המדידה (למשל, אתחלתם ב-CPU ולכן הם ב-RAM, וה-kernel חייב להגר אותם ל-HBM3).
תרגיל 4 - חפיפת העתקה וחישוב עם שני streams¶
- חזרו לניהול הידני (
cudaMallocעל הdevice) אבל הקצו את זיכרון הhosth_a,h_b,h_cכ-pinned עםcudaMallocHost. בחרוn = 1 << 25. - כתבו גרסה סדרתית:
cudaMemcpyH2D של שני הקלטים,vecAddעל כל המערך,cudaMemcpyD2H של הפלט. מדדו את זמנה הכולל עםcudaEvent. - כתבו גרסה חופפת: פרקו את המערך ל-
nChunks = 8נתחים, השתמשו בשני streams ב-round-robin, ולכל נתח הנפיקוcudaMemcpyAsyncH2D,vecAddעל הנתח, ו-cudaMemcpyAsyncD2H - הכל על ה-stream של הנתח. סנכרנו בסוף עםcudaStreamSynchronize. מדדו. - חשבו את ה-speedup בין הגרסה הסדרתית לחופפת. אמתו שהתוצאה זהה בשתי הגרסאות.
רמז: זכרו את שני התנאים ההכרחיים לחפיפה - זיכרון pinned והנפקה על streams לא-default. בלי pinned, cudaMemcpyAsync מסתנכרנת ולא תראו חפיפה כלל. הכתובות של הנתחים הן d_a + off ו-h_a + off כאשר off = c * chunk. חיבור וקטורים הוא memory-bound וזול בחישוב, ולכן החפיפה שתראו כאן היא בעיקר בין H2D ל-D2H של נתחים שונים; ה-speedup יהיה מתון (למשל 1.3x-1.7x). כדי לראות חפיפה דרמטית יותר, נסו בתרגיל הבונוס kernel כבד יותר.
תרגיל 5 - תזמון kernel מדויק עם cudaEvent¶
- תזמנו רק את ה-
vecAdd(בלי ההעברות) עבורn = 1 << 27עםcudaEvent, אחרי איטרציית חימום. הריצו 20 פעמים ודווחו חציון. - חשבו את רוחב הפס האפקטיבי של ה-kernel: כל איבר קורא שני
floatוכותב אחד, כלומר12 * nבתים; רוחב הפס הוא(12 * n) / (ms * 1e6)GB/s. - השוו לרוחב הפס התיאורטי של הכרטיס (H100 SXM: כ-3350 GB/s). איזה אחוז ניצלתם?
- מדדו בנפרד גם את זמן ה-H2D של שני הקלטים, ודווחו איזה מהם - ה-kernel או ההעברות - שולט בזמן הכולל.
רמז: cudaEventElapsedTime מחזיר מילישניות כ-float ברזולוציה של כחצי מיקרו-שנייה. את החימום עושים כדי לא למדוד את ה-JIT של PTX ל-SASS ואת ההקצאות הראשונות. תגלו שזמן ה-kernel זעום ורוחב הפס גבוה (עשרות אחוזים מהתיאורטי - סימן ל-coalescing טוב), אבל זמן ה-H2D על ה-PCIe גדול ממנו בהרבה. זו בדיוק הסיבה שלמדנו pinned וחפיפה.
תרגיל 6 (בונוס) - oversubscription וקרנל כבד¶
- Oversubscription: נסו
cudaMallocManagedעל נפח גדול מזיכרון ה-GPU (למשל אם יש לכם 24 GiB, הקצו שלושה מערכים של 10 GiB כל אחד). ודאו שההקצאה מצליחה (בניגוד ל-cudaMallocשהיה נכשל ב-out of memory), הריצו kernel פשוט שנוגע בכל האיברים, ומדדו כמה איטי זה בגלל ההחלפה. - קרנל כבד לחפיפה טובה יותר: החליפו את
vecAddב-kernel שמבצע הרבה חישוב לכל איבר (למשל לולאתfor (int k = 0; k < 256; k++) v = v*1.0001f + 0.5f;), חזרו על תרגיל 4, וראו איך ה-speedup מהחפיפה גדל כשזמן החישוב מתקרב לזמן ההעתקה. - הסבירו לפי מודל ה-
max(copy, compute)מההרצאה למה קרנל כבד יותר מפיק חפיפה טובה יותר, ומתי החפיפה מפסיקה לעזור (כשאחד הרכיבים שולט לגמרי).
רמז: ל-oversubscription, השתמשו ב-cudaMemGetInfo(&freeB, &totalB) כדי לדעת כמה זיכרון יש, ואז הקצו יותר. לחפיפה המשופרת, ככל שהחישוב לנתח מתקרב לזמן ההעתקה שלו כך max(copy, compute) קרוב ל-copy + compute חלקי 2 - החיסכון מקסימלי. אם החישוב שולט לגמרי (קרנל כבד מאוד), ההעתקות "מתחבאות" מתחתיו וזמן הריצה מתכנס לזמן החישוב לבדו.