3.2 ניהול זיכרון מהמארח הרצאה

בשיעור 0.3 כתבנו את חיבור הווקטורים הראשון וראינו את כוריאוגרפיית חמשת השלבים - cudaMalloc, cudaMemcpy H2D, launch, cudaMemcpy D2H, cudaFree - ובשיעור 3.1 מיפינו את היררכיית הזיכרון של הdevice מבפנים (registers, shared memory, global memory) מנקודת המבט של ה-kernel. בשיעור הזה אנחנו חוזרים אל צד הhost (כלומר ה-CPU) ושואלים את השאלה המעשית שמלווה כל תוכנית CUDA רצינית: איך מנהלים את הזיכרון של הdevice ביעילות מהhost? נעמיק בכל אחד מכלי הניהול - ההקצאה וההעתקה הבסיסיות, זיכרון host נעול-עמודים (pinned) שפותח את הדלת להעברות מהירות ואסינכרוניות, זיכרון מנוהל (managed / unified) שמטשטש את הגבול בין המרחבים ומעביר עמודים לפי דרישה, וה-streams (זרמים) שמאפשרים לחפוף העתקה עם חישוב במקום לשלם עליהם בזה אחר זה. את ה-streams נציג כאן ברמת המבוא ההכרחית כדי לחפוף העברות, ואת הטיפול המלא בהם - תלויות, עדיפויות, callbacks - נפרוש בשיעור 5.3. המסר המרכזי: לרוב, צוואר הבקבוק של תוכנית CUDA איננו ה-kernel אלא תנועת הנתונים סביבו, וכאן נלמד לשלוט בה.

הבסיס - ההקצאה, ההעתקה והכיוונים - allocate, copy, directions¶

נחזור בקצרה על שלישיית הפונקציות מ-CUDA Runtime API (הinterface ברמה הגבוהה עם התחילית cuda, כפי שראינו בשיעור 0.4), הפעם עם דגש על הפרטים שדילגנו עליהם. cudaMalloc(void** ptr, size_t bytes) מקצה בלוק רציף ב-global memory של הdevice ומחזיר קוד שגיאה; הכתובת נכתבת לתוך המצביע שלכם. cudaFree(ptr) משחרר אותו. cudaMemcpy(dst, src, bytes, kind) מעתיק בתים בין המרחבים, כשהפרמטר kind קובע את הכיוון:

ערך `kind`	כינוי	כיוון
`cudaMemcpyHostToDevice`	H2D	host -> device (מעלים קלט)
`cudaMemcpyDeviceToHost`	D2H	device -> host (מורידים תוצאה)
`cudaMemcpyDeviceToDevice`	D2D	device -> device (בתוך אותו GPU)
`cudaMemcpyHostToHost`	H2H	host -> host (נדיר, כמו `memcpy`)

העתקת D2D שימושית כשמעתיקים בין שני buffers על אותו כרטיס בלי לצאת אל הhost - היא רצה על רוחב הפס של ה-HBM3 (ב-H100, בסדר גודל של 3.35 TB/s) ולא על ה-PCIe האיטי. יש גם cudaMemcpyDefault, שמסיק את הכיוון לבד מסוג המצביעים באמצעות Unified Virtual Addressing - נוח, אבל עדיף לציין כיוון מפורש כדי שהקוד יתעד את כוונתו.

הבחנה שכדאי לדעת: כל הפונקציות האלה שייכות ל-CUDA Runtime API. יש להן מקבילות ברמה נמוכה יותר ב-CUDA Driver API (עם התחילית cu במקום cuda, למשל cuMemAlloc ו-cuMemcpyHtoD). ה-Runtime API הוא עטיפה מעל ה-Driver API, וכמעט כל התוכניות משתמשות ב-Runtime בגלל הארגונומיה; את ה-Driver API שומרים לשליטה מפורשת ב-contexts ובהשקות. בקורס הזה נשארים ב-Runtime API לכל אורך הדרך.

נקודה קריטית שנחזור אליה מיד: cudaMemcpy הרגילה היא חוסמת (blocking / synchronous) מנקודת מבט הhost - היא לא חוזרת עד שההעתקה הסתיימה. זה נוח וקל להבנה, אבל זה גם אומר שבזמן ההעברה ה-CPU עומד ומחכה, וה-GPU אינו יכול לחפוף אליה חישוב. כדי לשבור את המחסום הזה נצטרך שני דברים חדשים: זיכרון host מסוג pinned, ו-copy אסינכרוני על stream. נתחיל מהזיכרון.

זיכרון נעול-עמודים - pinned (page-locked) memory¶

כדי להבין למה זיכרון pinned מהיר יותר, צריך להיזכר איך מערכת ההפעלה מנהלת את זיכרון הhost. הקצאה רגילה ב-malloc מחזירה זיכרון בר-החלפה (pageable): מערכת ההפעלה רשאית להזיז את העמודים הפיזיים שלו, אפילו לגרש אותם אל הדיסק (swap), ולכן הכתובת הפיזית של buffer יכולה להשתנות מתחת לרגליים. ה-DMA engine (מנוע ההעברה הישירה) שעל ה-GPU, לעומת זאת, עובד מול כתובות פיזיות וחייב שהן יישארו קבועות לאורך כל ההעברה.

מכאן נובע מנגנון ה-bounce buffer (חוצץ ביניים) שקורה מאחורי הקלעים כשמעתיקים מזיכרון pageable:

   cudaMemcpy from pageable memory (the slow path):

   [ h_a - pageable ]                    [ d_a - GPU ]
        │                                     ▲
        │ 1. the driver copies (CPU memcpy)     │ 2. DMA
        ▼    to an internal pinned buffer         │    from the pinned buffer
   [ pinned staging buffer ]  ──────────────────┘

   two copies: first CPU-copy to the buffer, then DMA to the GPU.

הדרייבר מקצה חוצץ נעול פנימי, מעתיק אליו את הנתונים בעותק CPU רגיל, ורק ממנו מבצע את ה-DMA אל הdevice. עותק כפול זה מבזבז רוחב פס ומחייב את ה-CPU. אם נקצה את זיכרון הhost מלכתחילה כנעול-עמודים (page-locked, המכונה pinned), מערכת ההפעלה מבטיחה שהעמודים לא יזוזו ולא יגורשו, כך שה-DMA engine יכול לקרוא ישירות מזיכרון הhost - העתקה אחת בלבד:

   cudaMemcpy from pinned memory (the fast path):

   [ h_a - pinned ]  ───── DMA ─────►  [ d_a - GPU ]
   a single direct copy, no CPU involvement.

שתי דרכים להקצות זיכרון pinned, ובן זוגן לשחרור:

float* h_a;
CUDA_CHECK(cudaMallocHost(&h_a, bytes));      // the simple way
// ... use ...
CUDA_CHECK(cudaFreeHost(h_a));

// a richer version with flags:
float* h_b;
CUDA_CHECK(cudaHostAlloc(&h_b, bytes, cudaHostAllocDefault));
CUDA_CHECK(cudaFreeHost(h_b));

הפונקציה cudaMallocHost היא הקיצור הנוח. cudaHostAlloc חושפת דגלים מתקדמים: cudaHostAllocPortable (הזיכרון נעול עבור כל ה-contexts), cudaHostAllocMapped (ממפה את הזיכרון גם למרחב הכתובות של הdevice - zero-copy), ו-cudaHostAllocWriteCombined (cache write-combining, מהיר יותר להעברות H2D אבל איטי מאוד לקריאה מה-CPU). לרוב המקרים cudaMallocHost מספיקה. חובה לשחרר זיכרון pinned עם cudaFreeHost ולא עם free - אלה מקצים ממאגרים שונים לגמרי.

שני היתרונות של pinned, וזו הנקודה המרכזית של הסעיף:

מהירות H2D/D2H גבוהה יותר - העתקה אחת במקום שתיים. בפועל, על מערכת עם PCIe Gen5, העברה מזיכרון pinned מגיעה קרוב לרוחב הפס התיאורטי של ה-PCIe (כ-64 GB/s, בפועל 50-55 GB/s), בעוד שהעברה מזיכרון pageable נעצרת לרוב סביב חצי מכך. פי 2 עד פי 3 הבדל הוא תוצאה טיפוסית שתמדדו בעצמכם בתרגול.
תנאי הכרחי להעברות אסינכרוניות - רק העתקה מזיכרון pinned יכולה להיות באמת אסינכרונית ולחפוף לחישוב. את זה נפתח בסעיף ה-streams.

מחיר ה-over-pinning (מלכוד). זיכרון pinned אינו קסם חינם. עמודים נעולים אינם ניתנים להחלפה, ולכן הם גוזלים באופן קשיח מזיכרון ה-RAM הפיזי הזמין למערכת ההפעלה ולשאר התהליכים. אם תנעלו נתחים ענקיים - נניח רוב ה-RAM - תחניקו את המערכת כולה: תהליכים אחרים ייכנסו ל-swap, וההקצאה עצמה תואט (נעילת עמודים היא קריאת מערכת יקרה). לכן הכלל: נעלו רק את ה-buffers שבאמת חוזרים ומעורבים בהעברות (למשל buffers של קלט/פלט בלולאה חמה), לא כל הקצאת host. הקצאת pinned גם איטית יותר מ-malloc, ולכן מקצים אותה פעם אחת ומחזירים אותה שימוש, לא בכל איטרציה.

זיכרון מנוהל - unified / managed memory¶

עד עכשיו ניהלנו שני מרחבי זיכרון נפרדים ידנית: מצביע host, מצביע device, והעתקות מפורשות ביניהם. זיכרון מנוהל (managed memory, המכונה גם unified memory) מציע מודל אחר: הקצאה אחת שמצביע יחיד עליה תקף גם על הhost וגם על הdevice, וה-CUDA runtime מזיז את הנתונים ביניהם אוטומטית לפי הצורך.

float* a;
CUDA_CHECK(cudaMallocManaged(&a, bytes));   // single pointer, both worlds
for (int i = 0; i < n; i++) a[i] = 1.0f;    // write from the host - legal!
vecAdd<<<blocks, threads>>>(a, b, c, n);    // same pointer to the kernel - legal!
CUDA_CHECK(cudaDeviceSynchronize());
printf("%f\n", c[0]);                        // read from the host - legal!
CUDA_CHECK(cudaFree(a));                      // yes, cudaFree (not cudaFreeHost)

שימו לב כמה זה קצר לעומת שרשרת חמשת השלבים. אין cudaMalloc נפרד לhost ולdevice, אין cudaMemcpy. איך זה עובד? על ארכיטקטורות Pascal ומעלה (וה-H100 בכללן), הזיכרון המנוהל נשען על migration של עמודים לפי דרישה (on-demand page migration) באמצעות page faults של החומרה:

   the sequence with managed memory:

   1. cudaMallocManaged  ─► logical allocation, not yet physically placed
   2. the host writes a[i]  ─► page fault -> the page is placed in host RAM
   3. the kernel reads a[i] ─► GPU page fault -> the page migrates to HBM3
   4. the host reads c[i]   ─► page fault -> the page migrates back to RAM

כשה-kernel ניגש לעמוד שנמצא כרגע בזיכרון הhost, ה-GPU מייצר page fault, ה-runtime עוצר, מעתיק את העמוד (בדרך כלל 4 KiB או 2 MiB) אל ה-HBM3, וממשיך. ההגירה אוטומטית לחלוטין. היתרון: פרודוקטיביות. מבנים מקוננים עם מצביעים, פורטינג מהיר של קוד CPU קיים, ופיתוח ראשוני - כל אלה נעשים פשוט הרבה יותר. יתרון שני: oversubscription (הקצאת-יתר). אפשר להקצות עם cudaMallocManaged יותר זיכרון ממה שנכנס פיזית ב-80 GiB של ה-HBM3; ה-runtime יחליף עמודים פנימה והחוצה לפי הצורך, בדיוק כמו זיכרון וירטואלי. עם cudaMalloc רגיל הקצאה כזו פשוט הייתה נכשלת ב-out of memory.

המחיר (caveat ביצועי). ה-page faults אינם חינם. אם ה-kernel נוגע לראשונה בכל הנתונים, כל אחד מהם מייצר fault, וההגירה ה"עצלה" עמוד-אחר-עמוד איטית בהרבה מ-cudaMemcpy אחד גדול ורציף. בקוד תמים, managed memory יכול להיות איטי משמעותית מניהול ידני, כי אתם משלמים על ההעברות בזמן הריצה של ה-kernel, מפוזרות ולא בבת אחת. הפתרון: לומר ל-runtime מראש לאן שייכים הנתונים.

Prefetch מפורש עם cudaMemPrefetchAsync. מזיזים בלוק שלם אל היעד לפני שנוגעים בו, במקום להמתין ל-faults:

int device = 0;
// before the launch: move a,b to the GPU at once (one large contiguous transfer)
CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, device));
CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, device));
vecAdd<<<blocks, threads>>>(a, b, c, n);
// after the launch: move the result c back to the host ahead of time
CUDA_CHECK(cudaMemPrefetchAsync(c, bytes, cudaCpuDeviceId));
CUDA_CHECK(cudaDeviceSynchronize());

ה-prefetch הופך את ההגירה מ"הרבה faults קטנים בזמן ה-kernel" ל"העברה אחת גדולה, מקונבעת, לפני ה-kernel" - וזה מחזיר את הביצועים קרוב לרמת הניהול הידני, בלי לוותר על נוחות המצביע היחיד. הקבוע cudaCpuDeviceId מציין את הhost כיעד.

רמזים עם cudaMemAdvise. כשהגישה לנתונים אינה חד-כיוונית, אפשר לרמוז ל-runtime על דפוס השימוש כדי שינהל את העמודים חכם יותר:

רמז (advice)	משמעות
`cudaMemAdviseSetReadMostly`	הנתונים נקראים בעיקר ולא נכתבים; ה-runtime רשאי לשכפל אותם (עותק לקריאה בכל הdevice)
`cudaMemAdviseSetPreferredLocation`	קבע "בית" מועדף לעמודים; הגירות ינסו להשאיר אותם שם
`cudaMemAdviseSetAccessedBy`	הdevice הזה ייגש לנתונים; מפה אותם למרחבו במקום להגר בכל fault

השימוש: cudaMemAdvise(ptr, bytes, cudaMemAdviseSetPreferredLocation, device). הרמזים אינם משנים נכונות - רק ביצועים - ולכן הם כלי כוונון, לא חובה. הכלל המנחה: התחילו עם cudaMallocManaged לפרודוקטיביות, מדדו, ואם ה-page faults פוגעים - הוסיפו cudaMemPrefetchAsync בנקודות הברורות, ורק אם צריך עוד, כווננו עם cudaMemAdvise.

זרמים - streams - מבוא¶

עד עכשיו כל פעולות ה-CUDA שלנו רצו זו אחר זו, סדרתית. stream (זרם) הוא תור של פעולות GPU המתבצעות לפי סדר. שתי עובדות היסוד:

בתוך אותו stream, הפעולות מבוצעות לפי הסדר שבו הונפקו. פעולה שנייה לא מתחילה עד שהראשונה הסתיימה.
בין streams שונים אין סדר מובטח - הם עצמאיים וה-GPU רשאי להריץ אותם במקביל, בכפוף למשאבי החומרה.

זהו בדיוק המנוף שמאפשר חפיפה (overlap): אם נשים העתקה של נתח אחד ב-stream אחד וחישוב על נתח אחר ב-stream שני, ה-GPU יכול לבצע את שניהם בו-זמנית, כי מנועי ה-DMA (Copy Engines) של ה-H100 עצמאיים מ-CUDA Cores. יצירה והשמדה של stream:

cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
// ... use ...
CUDA_CHECK(cudaStreamDestroy(stream));

הגרסה האסינכרונית של ההעתקה היא cudaMemcpyAsync, שמקבלת פרמטר stream נוסף וחוזרת לhost מיד, בלי לחכות לסיום:

CUDA_CHECK(cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream));
myKernel<<<blocks, threads, 0, stream>>>(d_a, ...);   // the fourth parameter is the stream
CUDA_CHECK(cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream));

שימו לב לתצורת הlaunch בעלת ארבעה שדות <<<grid, block, sharedBytes, stream>>> - השדה הרביעי הוא ה-stream (השלישי, כאן 0, הוא shared memory דינמי). תנאי הכרחי: כדי ש-cudaMemcpyAsync תהיה באמת אסינכרונית, זיכרון הhost חייב להיות pinned. אם תעבירו לה מצביע pageable, הקריאה תתנהג כמעט כמו cudaMemcpy חוסמת - הדרייבר נופל בחזרה ל-bounce buffer ולסנכרון. זו הסיבה ששני הסעיפים הקודמים מתחברים כאן: pinned + async + streams הם השלישייה שמאפשרת חפיפה.

מלכוד ה-default stream. אם לא תציינו stream (כמו בכל הקוד שלנו עד שיעור זה), הפעולות רצות ב-default stream (stream 0). ב-semantics המסורתי (legacy), ה-default stream מסונכרן מול כל שאר ה-streams - כלומר פעולה בו חוסמת את כולם, ולא תקבלו חפיפה. לכן, לחפיפה אמיתית, צרו streams מפורשים ושמו בהם את העבודה, ולעולם אל תשאירו את הפעולות הקריטיות ב-default stream.

מדידה ותלויות עם events - CUDA events¶

איך יודעים כמה זמן לקח ה-kernel, וכיצד יוצרים תלות בין streams? התשובה לשניהם היא event (אירוע) - סמן שמניחים בתוך ה-stream ושמוקלט ברגע שכל הפעולות שלפניו הסתיימו. ל-event שני תפקידים:

מדידת זמן מדויקת. בניגוד ל-std::chrono שמודד על ה-CPU (וממדוד גם overhead של הסנכרון), event מוקלט על ה-GPU עצמו ומודד את הזמן שבו ה-GPU באמת עבד:

cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));

CUDA_CHECK(cudaEventRecord(start));                 // mark before
myKernel<<<blocks, threads>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaEventRecord(stop));                  // mark after
CUDA_CHECK(cudaEventSynchronize(stop));             // wait for stop to be recorded
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop)); // difference in milliseconds (float)
printf("kernel: %.3f ms\n", ms);
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));

הדיוק של cudaEventElapsedTime הוא בסדר גודל של חצי מיקרו-שנייה - הרבה מעבר לרזולוציה שנחוצה לנו. חשוב: cudaEventRecord(start) בלי ציון stream מקליט ב-default stream, שמסנכרן מול השאר; אם אתם מתזמנים עבודה ב-stream ספציפי, העבירו אותו כפרמטר שני: cudaEventRecord(start, stream). ותמיד הריצו איטרציית חימום (warm-up) לפני המדידה - הקומפילציה של PTX ל-SASS (JIT) וההקצאות הראשונות מוסיפות overhead חד-פעמי שתזהם מדידה יחידה.

יצירת תלות בין streams. cudaStreamWaitEvent(streamB, event) גורם ל-streamB להמתין עד שה-event (שהוקלט אולי ב-streamA) מתרחש - כך בונים תלות בין זרמים בלי לחסום את הhost. את זה נרחיב בשיעור 5.3. אם ה-event משמש רק לסנכרון ולא למדידה, כדאי ליצור אותו עם cudaEventCreateWithFlags(&e, cudaEventDisableTiming) כדי לחסוך overhead.

חפיפה של העתקה וחישוב - overlapping copy with compute¶

נחבר הכל לתבנית שהיא הסיבה האמיתית שלמדנו pinned, async ו-streams. נניח שאנחנו מעבדים מערך גדול. במקום להעתיק את כולו ל-GPU, לחשב, ולהעתיק חזרה - שלושה שלבים סדרתיים - נפרק אותו ל-chunks (נתחים) ונשכן אותם ב-streams מתחלפים. בזמן שה-GPU מחשב על נתח אחד, ה-DMA engine מעתיק את הנתח הבא:

const int nStreams = 2;
cudaStream_t streams[nStreams];
for (int s = 0; s < nStreams; s++)
    CUDA_CHECK(cudaStreamCreate(&streams[s]));

int chunk = n / nChunks;                 // chunk size (assuming an even split)
size_t cbytes = (size_t)chunk * sizeof(float);

for (int c = 0; c < nChunks; c++) {
    int off = c * chunk;                 // chunk offset
    cudaStream_t st = streams[c % nStreams];   // round-robin across the streams
    CUDA_CHECK(cudaMemcpyAsync(d_a + off, h_a + off, cbytes,
                               cudaMemcpyHostToDevice, st));
    CUDA_CHECK(cudaMemcpyAsync(d_b + off, h_b + off, cbytes,
                               cudaMemcpyHostToDevice, st));
    vecAdd<<<(chunk + 255) / 256, 256, 0, st>>>(d_a + off, d_b + off,
                                                d_c + off, chunk);
    CUDA_CHECK(cudaMemcpyAsync(h_c + off, d_c + off, cbytes,
                               cudaMemcpyDeviceToHost, st));
}
for (int s = 0; s < nStreams; s++)
    CUDA_CHECK(cudaStreamSynchronize(streams[s]));   // wait for completion

הרעיון בציר הזמן. סדרתי לעומת חופף, כשזמן ההעתקה של נתח וזמן החישוב עליו דומים:

   sequential (single stream, no overlap):
   H2D₀ ─► K₀ ─► D2H₀ ─► H2D₁ ─► K₁ ─► D2H₁ ─► ...
   |<---------- everything one after another ----------->|

   overlapped (two streams, pinned):
   stream A:  H2D₀ ─► K₀ ─► D2H₀ ─► H2D₂ ─► K₂ ─► D2H₂
   stream B:        H2D₁ ─► K₁ ─► D2H₁ ─► H2D₃ ─► K₃
                    ▲ while A computes, B copies - in parallel

חשבון מודל. נניח ארבעה נתחים, וזמן העתקה של נתח (H2D בלבד) שווה לזמן החישוב עליו, כל אחד t. סדרתי לחלוטין (H2D ואז K לכל נתח, בלי חפיפה) עולה בערך 2t x 4 = 8t. עם חפיפה מלאה בין ההעתקות לחישוב, הזמן מתקרב ל-max(סך ההעתקות, סך החישוב) ועוד עלות של נתח בודד בקצוות - בערך 4t + t = 5t. חיסכון של כ-37%. ככל שמגדילים את מספר הנתחים, "הזנב" של נתח בודד נעשה זניח, ומגיעים לחסם התיאורטי max(copy, compute) - הזמן של הרכיב האיטי לבדו, כאילו השני חינם. את זה בדיוק תמדדו בתרגול.

שני תנאים הכרחיים לחפיפה שקל לשכוח: (1) זיכרון הhost h_a, h_b, h_c חייב להיות pinned, אחרת cudaMemcpyAsync מסתנכרנת ואין חפיפה; (2) העבודה חייבת להתפזר על streams לא-default, אחרת ה-default stream מסדרן הכל. ה-H100 מצויד במספר Copy Engines, כך שהוא יכול אפילו לחפוף H2D ו-D2H בו-זמנית עם חישוב - שלושה זרמי עבודה במקביל.

מתי להשתמש במה - a decision table¶

נסכם את ארבעת הכלים כטבלת החלטה מעשית:

צורך	הכלי	הערה
הקצאה בסיסית של global memory	`cudaMalloc` / `cudaFree`	ברירת המחדל; שליטה מלאה, ניהול ידני
קלט/פלט חוזרים בהעברות	`cudaMallocHost` (pinned)	פי 2-3 מהירות; אל תנעלו יותר מדי
חפיפת העתקה עם חישוב	pinned + `cudaMemcpyAsync` + streams	דורש את שלושתם יחד
פרודוקטיביות, פורטינג, מבני מצביעים	`cudaMallocManaged`	הוסיפו `cudaMemPrefetchAsync` לביצועים
הקצאה גדולה מ-80 GiB (oversubscription)	`cudaMallocManaged`	ה-runtime מחליף עמודים אוטומטית
מדידת זמן GPU מדויקת	`cudaEvent`	תמיד עם איטרציית חימום

הכלל המנחה לאורך הקורס: צוואר הבקבוק לרוב הוא תנועת הנתונים, לא ה-kernel. ניהול נכון של הזיכרון מהhost - מזעור העברות, נעילת ה-buffers הנכונים, וחפיפת מה שנשאר עם חישוב - הוא לעיתים קרובות מה שמפריד בין קוד CUDA איטי לקוד CUDA מהיר, יותר מכל אופטימיזציה בתוך ה-kernel עצמו.

סיכום¶

cudaMemcpy הרגילה חוסמת את הhost; הכיוונים הם H2D, D2H, D2D (על רוחב הפס של ה-HBM3) ו-H2H, וכולם שייכים ל-CUDA Runtime API שהוא עטיפה מעל ה-Driver API.
זיכרון host pinned (נעול-עמודים) שמוקצה ב-cudaMallocHost או cudaHostAlloc מדלג על ה-bounce buffer ומאפשר DMA ישיר, ומכאן פי 2-3 מהירות בהעברות והתנאי ההכרחי להעברות אסינכרוניות.
נעילת-יתר (over-pinning) גוזלת RAM פיזי קשיח ומחניקה את המערכת; נועלים רק את ה-buffers החוזרים בהעברות, ומקצים אותם פעם אחת.
cudaMallocManaged מספק מצביע יחיד לשני העולמות עם migration של עמודים לפי דרישה (page faults); הוא מנצח בפרודוקטיביות וב-oversubscription אך משלם ב-page faults מפוזרים.
cudaMemPrefetchAsync ממיר את ההגירה מ-faults עצלים להעברה אחת גדולה ומחזיר ביצועים, ו-cudaMemAdvise מרמז על דפוסי גישה (ReadMostly, PreferredLocation, AccessedBy) לכוונון בלבד.
stream הוא תור פעולות מסודר; בתוך stream יש סדר, בין streams שונים אין - וזה בדיוק מה שמאפשר חפיפה.
cudaMemcpyAsync על stream חוזרת מיד, אך רק מזיכרון pinned ורק ב-stream לא-default היא באמת אסינכרונית וחופפת.
cudaEvent (עם cudaEventRecord ו-cudaEventElapsedTime) מודד זמן GPU בדיוק תת-מיקרו-שנייה ובונה תלויות בין streams דרך cudaStreamWaitEvent; תמיד עם איטרציית חימום.
חפיפת העתקה עם חישוב על פני נתחים ושני streams מקרבת את הזמן מ-copy + compute אל max(copy, compute), ודורשת בו-זמנית pinned, async ו-streams לא-default.