0.3 התוכנית הראשונה חיבור וקטורים הרצאה

בשיעור 0.2 התקנו את סביבת העבודה, ראינו ש-nvcc מתקמפל, והרצנו deviceQuery כדי לוודא שה-GPU חי ומדבר איתנו. עד עכשיו רק הסתכלנו על החומרה מבחוץ. בשיעור הזה נכתוב את התוכנית האמיתית הראשונה שלנו ב-CUDA - חיבור וקטורים, ה-"Hello, World" של תכנות GPU - ונריץ אותה על החומרה. לאורך הדרך נפגוש את כל אבני הבניין שילוו אותנו בכל שיעור מכאן והלאה: פונקציית kernel המסומנת ב-__global__, ניהול זיכרון בין הhost (כלומר ה-CPU) לבין הdevice (כלומר ה-GPU), תצורת ההרצה בסוגריים המשולשים <<< >>>, חישוב האינדקס הקנוני של כל thread, ומקרו איתן לטיפול בשגיאות. חיבור וקטורים הוא דוגמה פשוטה בכוונה - היא לא תנצל אפילו אחוז מכוח ה-GPU - אבל היא חושפת את כל השלד של תוכנית CUDA, ועליו נבנה את כל הקורס. את המודל התיאורטי מאחורי ה-threads וה-blocks נעמיק בפרק 2, ואת היררכיית הזיכרון בפרק 3; כאן המטרה היא לגעת בקוד רץ.

מהו kernel - the kernel¶

ה-kernel (קרנל, פונקציית ה-GPU) הוא יחידת הקוד שאנחנו כותבים ומרכיבים ב-CUDA, בדומה לפונקציה בשפה שרצה על CPU. אבל יש הבדל אחד עמוק ומגדיר שחייבים להפנים כבר עכשיו: kernel נקרא (launched) פעם אחת, אבל מבוצע פעמים רבות - פעם אחת בידי כל thread. בעולם ה-CPU קריאה אחת לפונקציה שווה הרצה אחת של גופה. ב-CUDA, launch אחת של kernel יוצרת אלפי הרצות מקבילות של אותו גוף קוד בדיוק, כל אחת בידי thread אחר.

מכאן נגזרות שלוש עובדות שמעצבות את כל האופן שבו חושבים על kernel:

kernel מחזיר תמיד void. הוא לא מחזיר ערך. התוצר היחיד שלו הוא תופעת הלוואי - כתיבות לזיכרון. הניסוח שנצמד אליו לאורך הקורס: kernel "מקבל מצביעים, ומשנה זיכרון" (just mutates memory). אם אתם מצפים ש-kernel "יחזיר" את התוצאה כמו פונקציה רגילה - זו התפיסה השגויה הנפוצה ביותר של מתחילים.
kernel מקבל מצביעים לזיכרון הגלובלי - global memory - של הdevice. הhost מעביר ל-kernel כתובות של זיכרון שנמצא על ה-GPU, וה-kernel קורא וכותב דרכן.
ההרצות בדרך כלל מקביליות וקונקורנטיות, בסדר לא דטרמיניסטי. אסור להניח שום סדר בין ה-threads.

הנה ה-kernel המלא לחיבור וקטורים. שימו לב לסימון __global__, לטיפוס ההחזרה void, ולכך שכל הפרמטרים הם מצביעים (או סקלר פשוט):

__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // the global index of the thread
    if (i < n) {                                    // bounds guard
        c[i] = a[i] + b[i];
    }
}

הערה חשובה על שמות: המילה __global__ אין לה שום קשר ל-global memory. היא מסמנת פונקציה שמושקת מהhost ומורצת על הdevice. אחד ממעצבי CUDA המוקדמים אמר בבדיחות הדעת שהשם נבחר "כדי למקסם בלבול אצל המפתחים". זכרו את ההבחנה: __global__ = kernel; global memory = הזיכרון על ה-GPU. אלה שני דברים שונים לגמרי שחולקים מילה במקרה.

שני עולמות - host וdevice - host and device¶

תוכנית CUDA חיה בשני עולמות בו-זמנית. יש קוד שרץ על הhost (ה-CPU), ויש קוד שרץ על הdevice (ה-GPU), ולכל אחד מהם מרחב זיכרון נפרד משלו. ה-CPU רואה את זיכרון ה-RAM הרגיל של המחשב; ה-GPU רואה את ה-HBM3 שלו (ב-H100 - 80 GiB של global memory). שני מרחבי הזיכרון האלה נפרדים פיזית ומחוברים בערוץ איטי יחסית (PCIe או NVLink).

מכאן נובע הכלל הקריטי ביותר של השיעור: מצביע לdevice (device pointer) אי אפשר לפענח (dereference) על הhost. אם cudaMalloc החזיר לכם כתובת של זיכרון על ה-GPU, וכתבתם על ה-CPU d_a[0] = 5; - התוכנית תקרוס (segmentation fault) או, גרוע מכך, תכתוב לכתובת אקראית ב-RAM של הhost. הכתובת שקיבלתם תקפה רק בתוך מרחב הכתובות של ה-GPU. באותו אופן, מצביע host רגיל אסור לפענח בתוך kernel.

לכן זרימת החיים של כמעט כל תוכנית CUDA היא אותה כוריאוגרפיה בת חמישה שלבים:

        Host (CPU / RAM)                             Device (GPU / HBM3)
   ┌───────────────────────┐                   ┌───────────────────────┐
   │ h_a, h_b  (input)      │                   │                       │
   │ h_c       (output)     │                   │  d_a, d_b, d_c        │
   └───────────────────────┘                   └───────────────────────┘
        1. cudaMalloc  ─────────────────────────►  allocate d_a,d_b,d_c
        2. cudaMemcpy(H2D) ─────────────────────►  h_a,h_b  →  d_a,d_b
        3. vecAdd<<<...>>>(d_a,d_b,d_c,n)  ──────►  kernel runs, writes to d_c
        4. cudaMemcpy(D2H) ◄─────────────────────  d_c  →  h_c
        5. cudaFree(d_a,d_b,d_c)  ──────────────►  free device memory

נעבור על ה-API של ניהול הזיכרון. שלוש הפונקציות הן חלק מ-CUDA Runtime API (הinterface ברמה הגבוהה שרוב התוכניות משתמשות בו, עם התחילית cuda):

cudaMalloc(void** ptr, size_t bytes) - מקצה bytes בתים ב-global memory של הdevice. שימו לב: הפרמטר הראשון הוא void** - כתובת של מצביע - כי הפונקציה כותבת לתוך המצביע שלכם את הכתובת החדשה על ה-GPU. הערך המוחזר של הפונקציה הוא קוד שגיאה, לא הכתובת.
cudaMemcpy(dst, src, bytes, kind) - מעתיק בתים בין המרחבים. הפרמטר kind קובע את הכיוון: cudaMemcpyHostToDevice (מכונה H2D, מעלים קלט ל-GPU) או cudaMemcpyDeviceToHost (מכונה D2H, מורידים תוצאה חזרה). יש גם cudaMemcpyDeviceToDevice.
cudaFree(ptr) - משחרר הקצאת device. בדיוק כמו free על ה-CPU, דליפת זיכרון על ה-GPU היא באג אמיתי.

הנה החלק הזה של הקוד בשלמותו:

int n = 1 << 20;                       // 1,048,576 elements
size_t bytes = (size_t)n * sizeof(float);   // 4 MiB per vector

float *h_a = (float*)malloc(bytes);    // regular host allocations
float *h_b = (float*)malloc(bytes);
float *h_c = (float*)malloc(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;                // device pointers
cudaMalloc(&d_a, bytes);               // pass &d_a since the function writes into it
cudaMalloc(&d_b, bytes);
cudaMalloc(&d_c, bytes);

cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);   // H2D
cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

אינדוקס וlaunch - launch configuration and indexing¶

הגענו ללב הרעיון של CUDA. השקנו kernel אחד שירוץ באלפי עותקים; איך כל עותק יודע על איזה איבר במערך לעבוד? התשובה היא אינדוקס מבוסס מזהים. כל thread מקבל שני מזהים מובנים שהחומרה ממלאת עבורו, ומהם הוא מחשב את מיקומו הגלובלי.

תחילה הlaunch עצמה. הסינטקס <<<numBlocks, blockSize>>> בין שם ה-kernel לרשימת הארגומנטים הוא תצורת ההרצה (execution configuration). הוא אומר לחומרה: הרץ את ה-kernel על grid (רשת) של numBlocks בלוקים, כשכל block מכיל blockSize threads:

int blockSize = 256;                              // threads per block
int numBlocks = (n + blockSize - 1) / blockSize;  // number of blocks
vecAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

בתוך ה-kernel, כל thread רואה ארבעה משתנים מובנים:

משתנה	משמעות
`threadIdx.x`	האינדקס של ה-thread בתוך ה-block שלו (0 עד `blockDim.x - 1`)
`blockDim.x`	מספר הthreads ב-block (כאן 256)
`blockIdx.x`	האינדקס של ה-block בתוך ה-grid
`gridDim.x`	מספר הבלוקים ב-grid

מהם נגזר האינדקס הגלובלי הקנוני - הביטוי שתראו בכל תוכנית CUDA:

int i = blockIdx.x * blockDim.x + threadIdx.x;

בואו נתרגם אותו במספרים. נניח blockDim.x = 256. ה-thread הראשון של block מספר 0 מקבל 0*256 + 0 = 0. ה-thread האחרון של block מספר 0 מקבל 0*256 + 255 = 255. ה-thread הראשון של block מספר 1 מקבל 1*256 + 0 = 256. וכך הלאה - כל האיברים 0..n-1 מכוסים בדיוק פעם אחת, בלי חפיפה ובלי חורים:

 block 0                block 1                block 2
 ┌─────────────────┐    ┌─────────────────┐    ┌─────────────────┐
 │ tid 0 .. 255    │    │ tid 0 .. 255    │    │ tid 0 .. 255    │
 └─────────────────┘    └─────────────────┘    └─────────────────┘
   i = 0 .. 255           i = 256 .. 511         i = 512 .. 767   ...

חישוב גודל ה-grid. אנחנו רוצים לפחות thread אחד לכל איבר. אם נחלק n / blockSize בחלוקה שלמה, ונשאר שארית, נפספס את האיברים האחרונים. לכן משתמשים בעיגול כלפי מעלה בנוסחה הסטנדרטית:

int numBlocks = (n + blockSize - 1) / blockSize;

עבור n = 1048576 ו-blockSize = 256 יוצא (1048576 + 255) / 256 = 4096 בלוקים בדיוק (כאן החלוקה יפה). אבל נניח n = 1000: אז (1000 + 255) / 256 = 4 בלוקים, כלומר 4 * 256 = 1024 threads - 24 threads יותר מהצורך. הthreads העודפים (i = 1000..1023) הם בדיוק הסיבה לשומר הגבולות (bounds guard):

if (i < n) { c[i] = a[i] + b[i]; }

בלי ה-if, אותם 24 threads היו קוראים וכותבים מעבר לסוף המערכים - גישה מחוץ לתחום שתגרום לשגיאה או להשחתת זיכרון. שומר הגבולות הוא לא קישוט; הוא הכרחי בכל פעם שה-grid לא מחלק בדיוק את n.

למה blockSize כפולה של 32? בחומרה, הthreads מתבצעים בקבוצות של 32 הנקראות warp (קבוצת 32 threads המבצעים אותה פקודה יחד). את המנגנון הזה נלמד לעומק בפרק 2, אבל כלל האצבע כבר עכשיו: בחרו blockSize שהיא כפולה של 32, אחרת ה-warp האחרון בכל block יהיה חלקי וחלק מנתיבי החישוב יתבזבזו. ערכים נפוצים הם 128, 256 ו-512; 256 הוא ברירת מחדל טובה. ה-H100 מאפשר עד 1024 threads ב-block, ועד 2048 threads תושבים (resident) בו-זמנית על כל אחד מ-132 ה-SMs שלו.

התוכנית המלאה - the complete program¶

נחבר את הכל לקובץ אחד רץ, vecadd.cu. הוא כולל כבר את מקרו הבדיקה CUDA_CHECK שנסביר בסעיף הבא - קראו את הקוד קודם, ואז נפרק את הטיפול בשגיאות:

#include <cstdio>
#include <cstdlib>
#include <cmath>
#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 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;                            // 1,048,576
    size_t bytes = (size_t)n * sizeof(float);

    // allocation and init on the host
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) { h_a[i] = 1.0f; h_b[i] = 2.0f; }

    // allocation on the device
    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));

    // copy input host -> device
    CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    // kernel launch
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vecAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
    CUDA_CHECK(cudaGetLastError());       // catches launch errors
    CUDA_CHECK(cudaDeviceSynchronize());  // catches runtime errors

    // copy result device -> host
    CUDA_CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

    // validation
    double maxErr = 0.0;
    for (int i = 0; i < n; i++) {
        maxErr = fmax(maxErr, fabs((double)h_c[i] - 3.0));
    }
    printf("n = %d, max error = %f\n", n, maxErr);
    printf("c[0] = %.1f, c[%d] = %.1f\n", h_c[0], n - 1, h_c[n - 1]);

    // free
    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));
    free(h_a); free(h_b); free(h_c);
    return 0;
}

הקומפילציה וההרצה, עם ה-H100 כמטרה (Hopper הוא sm_90a):

nvcc -O2 -arch=sm_90a -o vecadd vecadd.cu
./vecadd

הפלט הצפוי:

n = 1048576, max error = 0.000000
c[0] = 3.0, c[1048575] = 3.0

השגיאה המקסימלית היא 0 כי 1.0f + 2.0f == 3.0f בדיוק בייצוג הנקודה הצפה. אם אין לכם GPU מקומי, כל הקוד רץ ללא שינוי ב-Google Colab (עם T4 חינמי) או על מכונת ענן שכורה - כפי שראינו בשיעור 0.2. אם השמטתם את -arch, ה-nvcc יתקמפל לארכיטקטורת ברירת מחדל ישנה יותר, וזה עדיין ירוץ, אבל ציון הארכיטקטורה הנכונה מפיק SASS (אסמבלר זורם, קוד המכונה של ה-GPU) מותאם ל-H100.

טיפול בשגיאות - error handling¶

בקוד למעלה עטפנו כמעט כל קריאת CUDA במקרו CUDA_CHECK. זה לא פדנטיות - זה הכרחי, כי שגיאות CUDA שקטות. רוב פונקציות ה-Runtime API מחזירות ערך מסוג cudaError_t. אם לא בודקים אותו, שגיאה פשוט נבלעת והתוכנית ממשיכה עם נתונים שגויים. המקרו לוכד את קוד השגיאה, מדפיס את הקובץ, השורה, הקריאה עצמה (#call) ואת ההודעה הקריאה מ-cudaGetErrorString, ואז יוצא.

אבל launch ה-kernel מיוחדת ומצריכה שתי בדיקות נפרדות. הסיבה עמוקה: הlaunch kernel היא אסינכרונית. השורה vecAdd<<<...>>>(...) חוזרת לhost כמעט מיד, לפני שה-kernel בכלל סיים (ולעיתים לפני שהתחיל) לרוץ על ה-GPU. לכן:

vecAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaGetLastError());       // (1) errors detected at launch time
CUDA_CHECK(cudaDeviceSynchronize());  // (2) errors detected at runtime

cudaGetLastError() תופס שגיאות שהתגלו כבר ברגע הlaunch - למשל תצורת הרצה לא חוקית (blockSize גדול מ-1024), שמחזירה cudaErrorInvalidConfiguration. שימו לב שאי אפשר לעטוף את הlaunch עצמה ב-CUDA_CHECK כי vecAdd<<<...>>>() אינה מחזירה cudaError_t - היא סינטקס מיוחד, ולכן שואלים בנפרד "מה הייתה השגיאה האחרונה?".
cudaDeviceSynchronize() חוסם את הhost עד שה-GPU מסיים, ומחזיר שגיאות שקרו במהלך הריצה עצמה - למשל גישה לא חוקית לזיכרון בתוך ה-kernel. בלי הסנכרון, שגיאת ריצה עלולה להתגלות רק מאוחר יותר, בקריאת CUDA אחרת לגמרי, ולבלבל אתכם לחלוטין.

עוד מלכוד קריטי: שגיאות CUDA "דביקות" (sticky). ברגע שקרתה שגיאה בזמן ריצה של kernel (כמו גישה מחוץ לתחום), ה-context של ה-GPU נכנס למצב שבור, וכל קריאת CUDA שלאחר מכן תחזיר את אותה שגיאה, גם אם היא כשלעצמה תקינה. cudaGetLastError גם מאפס את דגל השגיאה, כך שקריאה כפולה לו תחזיר cudaSuccess בפעם השנייה. מסיבה זו נהוג לקרוא ל-cudaGetLastError פעם אחת מיד אחרי הlaunch כדי "לנקות ולבדוק" בבת אחת. בתרגול תגרמו לשגיאה בכוונה כדי לראות בעיניים איך שני המנגנונים האלה מתנהגים.

לולאת grid-stride - the grid-stride loop¶

בגרסה שכתבנו, כל thread מטפל בדיוק באיבר אחד, ולכן חייבים להשיק לפחות n threads. זה עובד, אבל שביר: מה אם n גדול ממספר הthreads המקסימלי שנוח להשיק? מה אם רוצים לכוונן את מספר הבלוקים למספר ה-SMs של הכרטיס במקום ל-n? הפתרון הנקי, שהוא גם התבנית המקצועית המומלצת, הוא לולאת grid-stride:

__global__ void vecAddStride(const float* a, const float* b,
                             float* c, int n) {
    int idx    = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;   // total number of threads in the grid
    for (int i = idx; i < n; i += stride) {
        c[i] = a[i] + b[i];
    }
}

הרעיון: במקום שכל thread יעשה איבר אחד, כל thread מדלג על פני המערך בצעדים בגודל stride - המספר הכולל של הthreads ב-grid - ומטפל בכל האיברים שנופלים על מסלולו. אם השקנו 4096 threads ו-n = 1,000,000, ה-thread שמספרו הגלובלי 0 יטפל באיברים 0, 4096, 8192, ... וכן הלאה, עד שיכסה את חלקו. כל האיברים מכוסים, בכל גודל grid.

   n = 12,  stride = 4  (only four threads)
   thread 0:  i = 0, 4, 8
   thread 1:  i = 1, 5, 9
   thread 2:  i = 2, 6, 10
   thread 3:  i = 3, 7, 11

היתרונות של התבנית הזו:

גמישות בגודל. אותו kernel עובד לכל n, קטן או ענק, בלי לשנות את תצורת הlaunch.
ניתוק גודל ה-grid מגודל הבעיה. אפשר לבחור numBlocks שממקסם תפוסה (occupancy) על החומרה - למשל כפולה של מספר ה-SMs - במקום להיגרר אחרי n.
ניצול-חוזר של threads. thread שכבר "חם" בקובץ האוגרים ובcache ממשיך לעבוד, במקום לשלם על launch של עוד ועוד threads.
בדיקות תקינות. קל להריץ את אותו kernel עם grid קטן במכוון (אפילו block יחיד) כדי לבדוק נכונות, ואז להגדיל.

שימו לב שהתנאי i < n בלולאה משמש גם כשומר הגבולות - אין צורך ב-if נפרד. בתרגול תמירו את חיבור הווקטורים ללולאת grid-stride ותריצו אותה עם grid קטן מהנתונים כדי לראות שהיא עדיין נכונה.

מתי ה-GPU באמת מנצח - when the GPU wins¶

מפתה לחשוב שברגע שהעברנו חישוב ל-GPU, הוא יהיה מהיר יותר. עבור חיבור וקטורים זה פשוט לא נכון בקנה מידה קטן, ולמה חשוב להבין. חיבור וקטורים הוא חישוב memory-bound: לכל איבר אנחנו קוראים שני float וכותבים אחד - שלוש גישות זיכרון - ומבצעים חיבור בודד. הarithmetic intensity של פעולה אחת ל-12 בתים היא נמוכה מאוד. ה-GPU לא שם כדי לעשות חשבון; הוא שם כדי לזרום דרך המון נתונים. לכן:

בקנה מידה קטן (n של אלפים) הoverhead מנצחת: העלות של cudaMemcpy על גבי ה-PCIe ושל launch ה-kernel גדולה בהרבה מהחישוב עצמו, ולולאת CPU פשוטה תסיים לפני שבכלל העתקנו את הנתונים ל-GPU.
בקנה מידה גדול (n של עשרות ומאות מיליונים) זמן ה-kernel נשלט על ידי רוחב הפס של ה-HBM3 (ב-H100, בסדר גודל של 3.35 TB/s), שגבוה בהרבה מרוחב הפס של זיכרון ה-CPU. אז ה-GPU מנצח בבירור - כל עוד מודדים את זמן ה-kernel לבדו, או ששומרים את הנתונים על ה-GPU לאורך שרשרת שלמה של פעולות ולא משלמים על ההעברות בכל צעד.

זו הנקודה החינוכית: ההעברות על גבי ה-PCIe (בסדר גודל של עשרות GB/s) הן צוואר הבקבוק בחישוב כה קל. בתרגול תמדדו את שני העולמות עם cudaEvent ותראו במו עיניכם באיזה n הכף נוטה. את הכלים הפורמליים לניתוח - מודל ה-Roofline, arithmetic intensity ונקודת המפנה - נפתח בפרק על ביצועים.

סיכום¶

ה-kernel מסומן ב-__global__, מחזיר תמיד void, מקבל מצביעים ל-global memory, ו"רק משנה זיכרון" - הוא מושק פעם אחת ומבוצע בידי אלפי threads.
לhost (CPU) ולdevice (GPU) מרחבי זיכרון נפרדים; מצביע device אי אפשר לפענח על הhost, ולכן חייבים את שרשרת cudaMalloc -> cudaMemcpy H2D -> launch -> cudaMemcpy D2H -> cudaFree.
הפרמטר הראשון של cudaMalloc הוא void** כי הפונקציה כותבת את הכתובת לתוך המצביע שלכם; הערך המוחזר הוא קוד שגיאה.
תצורת ההרצה <<<numBlocks, blockSize>>> קובעת את ה-grid, וכל thread מחשב את מיקומו הגלובלי לפי i = blockIdx.x * blockDim.x + threadIdx.x.
גודל ה-grid מחושב בעיגול כלפי מעלה (n + blockSize - 1) / blockSize, ותמיד נדרש שומר גבולות if (i < n) כדי להגן על הthreads העודפים.
בוחרים blockSize ככפולה של 32 (בדרך כלל 256) כדי שהthreads ייארזו לתוך warps שלמים.
הlaunch kernel היא אסינכרונית; בודקים cudaGetLastError() מיד אחריה (שגיאות launch) ואז cudaDeviceSynchronize() (שגיאות ריצה), ושגיאות CUDA "דביקות" - הן ממשיכות להופיע עד שמנקים אותן.
לולאת grid-stride (for (int i = idx; i < n; i += blockDim.x * gridDim.x)) מנתקת את גודל ה-grid מגודל הבעיה, נכונה לכל n, וניתנת להרצה עם grid קטן במכוון.
חיבור וקטורים הוא memory-bound; ה-GPU מנצח רק בקנה מידה גדול ורק כשלא משלמים על העברות PCIe בכל צעד.