לדלג לתוכן

2.4 סנכרון, cooperative groups ו warpgroups פתרון

פתרון - סנכרון, cooperative groups ו-warpgroups

נעבור תרגיל אחרי תרגיל. כל המספרים והזמנים כאן הם מ-H100 (SXM) עם CUDA 12.4; זמני הריצה המדויקים אצלכם ישתנו בין כרטיסים וגרסאות toolkit, וזה תקין - מה שחשוב הוא הכיוון והיחסים. מי שעובד על T4 יראה זמנים ארוכים יותר ולא יוכל להשתמש ב-__reduce_*_sync החומרתי, אבל כל שאר הקוד עובד. הקלט הוא מערך של 2^24 אחדוֹת, ולכן הסכום הנכון הוא בדיוק 16,777,216.

פתרון תרגיל 1 - reduction קלאסית עם shared memory

__global__ void reduceShared(const float* in, float* out, int n) {
    extern __shared__ float sdata[];
    int tid = threadIdx.x;
    int i   = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? in[i] : 0.0f;   // global -> shared
    __syncthreads();

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();                    // outside the if!
    }
    if (tid == 0) out[blockIdx.x] = sdata[0];
}

$ nvcc -arch=sm_90 reduce.cu -o reduce && ./reduce
GPU sum = 16777216.0,  reference = 16777216.0,  OK

לגבי סעיף 4: ל-block של 256 threads יש __syncthreads אחד אחרי הטעינה, ואז שלב לכל חזקה של 2 מ-128 עד 1, כלומר 128, 64, 32, 16, 8, 4, 2, 1 = 8 שלבים. סך הכל 9 מחסומים לכל block.

למה זה עבד: המחסום אחרי הטעינה מבטיח שכל ה-tile ב-shared memory לפני שמישהו קורא תא של שכן, והמחסום בסוף כל שלב צמצום מבטיח ששני החצאים מוכנים לפני השלב הבא. ה-__syncthreads נמצא מחוץ ל-if (tid < s) ולכן כל 256 ה-threads מגיעים אליו בכל איטרציה - חוק הזהב נשמר. איך להכליל: זהו דפוס ה-tree reduction הסטנדרטי; הוא עובד לכל block שהוא חזקה של 2, ולגדלים אחרים צריך ריפוד (padding) של ה-shared memory באפסים.

פתרון תרגיל 2 - reduction עם warp shuffles

__global__ void reduceShfl(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float val = (i < n) ? in[i] : 0.0f;

    val = warpReduceSum(val);               // each warp reduces, result in lane 0

    __shared__ float warpSums[32];
    int lane = threadIdx.x % warpSize;
    int wid  = threadIdx.x / warpSize;
    if (lane == 0) warpSums[wid] = val;
    __syncthreads();                        // the only barrier

    int numWarps = blockDim.x / warpSize;
    val = (threadIdx.x < numWarps) ? warpSums[lane] : 0.0f;
    if (wid == 0) val = warpReduceSum(val);
    if (threadIdx.x == 0) out[blockIdx.x] = val;
}

$ ./reduce      # after replacing the launched kernel with reduceShfl
GPU sum = 16777216.0,  reference = 16777216.0,  OK

לגבי הספירה (סעיף 4): כאן יש מחסום אחד בלבד (לעומת 9 בתרגיל 1), והוא נדרש רק כדי לוודא שכל 8 ה-warps כתבו את הסכום החלקי שלהם ל-warpSums לפני שה-warp הראשון קורא אותם. משתמשים ב-32 תאי shared בלבד (מערך warpSums), לעומת 256 תאים בתרגיל 1.

למה זה עבד: הצמצום התוך-warp רץ כולו באוגרים דרך __shfl_down_sync, בלי shared ובלי מחסום, כי 32 ה-threads של warp רצים ב-lockstep ממילא. רק הצירוף בין 8 ה-warps דורש shared memory ומחסום אחד. איך להכליל: כל reduction מדרגת שני מפלסים - קודם warp (זול, באוגרים), אז block (מחסום יחיד) - וזה בדיוק המבנה של ספריות כמו CUB. לגדלי block מעל 1024 threads ה-warpSums[32] מספיק כי המקסימום הוא 1024/32 = 32 warps.

פתרון תרגיל 3 - השוואת נכונות וביצועים

קוד התזמון:

    cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);
    reduceShfl<<<grid, block, 0>>>(d_in, d_partial, n);       // warmup
    CUDA_CHECK(cudaDeviceSynchronize());
    cudaEventRecord(a);
    for (int r = 0; r < 100; ++r)
        reduceShfl<<<grid, block, 0>>>(d_in, d_partial, n);
    cudaEventRecord(b);
    cudaEventSynchronize(b);
    float ms = 0; cudaEventElapsedTime(&ms, a, b);
    printf("reduceShfl: %.3f ms/launch\n", ms / 100.0f);

$ ./reduce
reduceShared: 0.121 ms/launch   OK
reduceShfl:   0.098 ms/launch   OK

שתי הגרסאות נכונות (OK). על H100, עבור 2^24 אלמנטים, גרסת ה-shuffle מהירה בכ-20% (0.098 מול 0.121 ms). על קלט גדול או block גדול ההפרש גדל.

למה זה עבד: גרסת ה-shared מבצעת 9 מחסומים ו-8 שלבי צמצום דרך ה-L1, ולכל שלב פוטנציאל ל-bank conflicts; גרסת ה-shuffle מבצעת מחסום אחד ומצמצמת באוגרים. ההפרש הוא בדיוק העלות של ה-__syncthreads החוזרים והגישות ל-shared. הערה: ב-reduction על אחדוֹת ה-kernel memory-bound - שתי הגרסאות קרובות לרוחב הפס של ה-HBM, ולכן ההפרש קטן יחסית; ב-reduction עם עבודה חישובית לכל אלמנט ההפרש בולט יותר. איך להכליל: תמיד דחפו את התיאום למטה בהיררכיה (warp לפני block); הרווח גדל ככל שיש יותר שלבי צמצום שאפשר להעביר לאוגרים.

פתרון תרגיל 4 - Cooperative Groups + atomic חוצה-blocks

__global__ void reduceCG(const float* in, float* out, int n) {
    cg::thread_block block = cg::this_thread_block();
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float val = (i < n) ? in[i] : 0.0f;
    val = cg::reduce(warp, val, cg::plus<float>());

    __shared__ float warpSums[32];
    if (warp.thread_rank() == 0)
        warpSums[warp.meta_group_rank()] = val;
    block.sync();

    if (warp.meta_group_rank() == 0) {
        val = (warp.thread_rank() < warp.meta_group_size()) ? warpSums[warp.thread_rank()] : 0.0f;
        val = cg::reduce(warp, val, cg::plus<float>());
        if (warp.thread_rank() == 0)
            atomicAdd(out, val);            // cross-block combine - the only channel
    }
}

הlaunch עם אתחול התא הגלובלי:

    float* d_total;
    CUDA_CHECK(cudaMalloc(&d_total, sizeof(float)));
    CUDA_CHECK(cudaMemset(d_total, 0, sizeof(float)));
    reduceCG<<<grid, block>>>(d_in, d_total, n);
    CUDA_CHECK(cudaGetLastError());
    float total = 0;
    CUDA_CHECK(cudaMemcpy(&total, d_total, sizeof(float), cudaMemcpyDeviceToHost));
    printf("total = %.1f\n", total);

$ ./reduce
total = 16777216.0    OK

התא הגלובלי היחיד מכיל את הסכום המלא, בלי לולאת סיכום על ה-CPU. תשובה לסעיף 4: השתמשנו ב-atomicAdd כי אין מחסום חוצה-blocks - ה-blocks יושבים על SMs שונים, מתוזמנים בכל סדר, וייתכן שחלקם עוד לא התחילו לרוץ כשאחרים סיימו. אילו ניסינו __syncthreads כדי לתאם blocks, זה כלל לא היה מסתנכרן ביניהם (__syncthreads פועל רק בתוך block), ואילו ניסינו busy-wait של block אחד על אחר, היינו מסתכנים ב-deadlock כי ה-block שמחכים לו אולי לעולם לא יתוזמן.

למה זה עבד: cg::reduce מחליף את לולאת ה-shuffle הידנית בקריאה אחת קריאה (ועל CC 8.0+ מתקמפל להוראת __reduce חומרתית), ו-atomicAdd מבצע את הצעד החוצה-blocks היחיד האפשרי - קריאה/כתיבה אטומית ל-global memory. איך להכליל: הדפוס "צמצם בתוך block, ואז atomicAdd אחד לתא גלובלי" הוא reduction בשלב יחיד; החיסרון הוא צוואר-בקבוק על ה-atomic היחיד כשיש הרבה blocks, ולכן ל-reduction בקנה מידה גדול משתמשים בשני kernels או ב-atomic מדורג - נרחיב בשיעור 9.4.

פתרון תרגיל 5 - צמצום ברמת warp עם tiled_partition

template<int TILE>
__device__ float tileReduce(cg::thread_block_tile<TILE> g, float v) {
    return cg::reduce(g, v, cg::plus<float>());
}

__global__ void tilesDemo(const float* in, float* out, int n) {
    cg::thread_block block = cg::this_thread_block();
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = (i < n) ? in[i] : 0.0f;

    auto warp32 = cg::tiled_partition<32>(block);
    float s32 = tileReduce(warp32, v);

    auto tile16 = cg::tiled_partition<16>(block);
    float s16 = tileReduce(tile16, v);

    if (tile16.thread_rank() == 0 && blockIdx.x == 0)
        printf("tile16 #%d: size=%d rank0-sum=%.0f  (warp32 sum=%.0f)\n",
               tile16.meta_group_rank(), tile16.size(), s16, s32);
}

$ ./reduce      # input of ones
tile16 #0: size=16 rank0-sum=16  (warp32 sum=32)
tile16 #1: size=16 rank0-sum=16  (warp32 sum=32)
tile16 #2: size=16 rank0-sum=16  (warp32 sum=32)
...

כל tile של 16 מצטמצם בנפרד (סכום 16 אחדוֹת = 16), ואילו warp של 32 נותן 32. meta_group_rank() מראה שכל warp מתחלק לשני tiles של 16.

למה זה עבד: tiled_partition<N> מחלק את ה-block לקבוצות עצמאיות בגודל N, כל אחת עם thread_rank, size ו-sync משלה; cg::reduce פועל על כל קבוצה בנפרד ומכבד את גבולותיה. איך להכליל: היתרון על פני __shfl_down_sync ידני הוא שהקבוצה נושאת את המסכה בתוכה - אין סיכון לטעות במסכת ה-lanes, והקוד גנרי (אותה פונקציה עובדת ל-32, 16, 8...). ה-tiles הקטנים מ-32 שימושיים כשכל קבוצה קטנה מטפלת ביחידת עבודה נפרדת, למשל 8 threads לכל שורה במטריצה קטנה או ב-softmax על וקטורים קצרים.

פתרון תרגיל 6 - למה מחסום מפוצל גורם ל-deadlock

ה-kernel השגוי:

__global__ void badSync(const float* in, float* out, int n) {
    __shared__ float s[256];
    int tid = threadIdx.x;
    if (tid < 100) {              // only 100 out of 256 enter
        s[tid] = in[tid];
        __syncthreads();          // deadlock! 156 threads never arrive
    }
    if (tid == 0) out[blockIdx.x] = s[0];
}

$ timeout 10 ./reduce
# ... no output, the process hangs for 10 seconds and then gets killed by timeout
$ echo $?
124        # timeout killed the process; the kernel never returned

אם היינו בודקים את cudaDeviceSynchronize, הוא לא היה חוזר כלל (או שהיה חוזר עם cudaErrorLaunchTimeout על מערכת עם watchdog). לגבי סעיף 3: המחסום ממתין שכל 256 ה-threads של ה-block יגיעו אליו, אבל רק 100 נכנסו לענף if (tid < 100) והגיעו למחסום. 156 ה-threads הנותרים דילגו על הענף כולו ולעולם לא יגיעו למחסום. 100 ה-threads שהגיעו ממתינים לנצח ל-156 שלא יבואו - קיפאון.

התיקון (סעיף 4):

__global__ void goodSync(const float* in, float* out, int n) {
    __shared__ float s[256];
    int tid = threadIdx.x;
    s[tid] = (tid < 100) ? in[tid] : 0.0f;  // the condition is in the value, not in control flow
    __syncthreads();                         // all 256 arrive
    if (tid == 0) out[blockIdx.x] = s[0];
}

למה זה עבד: __syncthreads הוא מחסום קולקטיבי שסופר את כל ה-threads של ה-block; אם ולו thread אחד לא מגיע אליו, ה-block קופא. איך להכליל: הכלל הוא "מחסום לעולם לא בתוך זרימת בקרה תלוית-thread". כל הגנה או תנאי שצריך - מבצעים בערך שנטען/נכתב (למשל טעינת 0.0f מחוץ לתחום), וה-__syncthreads עצמו נשאר בנתיב שכל ה-threads עוברים בו. אותו עיקרון חל על block.sync() של Cooperative Groups ועל כל מחסום קולקטיבי אחר.

פתרון תרגיל 7 (בונוס) - warpgroup ו-Tensor Cores

$ nvcc --list-gpu-arch
compute_50
...
compute_90
compute_90a      <- the a-suffix; required for wgmma

על H100 מופיע sm_90a/compute_90a, ו-wgmma.mma_async דורש אותו במפורש (-arch=sm_90a). על Colab עם T4 (CC 7.5) המושג warpgroup אינו קיים - הוא הוצג רק ב-Hopper, ו-T4 (Turing) אינו מכיר wgmma כלל.

תשובות ניתוח:

  • סעיף 2 - למה 4 warps: ראינו בפרק 1 ש-SM של H100 בנוי מארבע תת-מחיצות, כל אחת עם warp scheduler ו-Tensor Core משלה. warpgroup ממפה warp אחד לכל תת-מחיצה, כך שהוראת warpgroup מפעילה את כל ארבעת ה-Tensor Cores של ה-SM בו-זמנית. 3 warps היו משאירים תת-מחיצה בטלה; 8 warps לא היו נכנסים בארבע התת-מחיצות בלי כפילות. לכן דווקא 4, ולכן יישור ה-warp-rank הראשון לכפולה של 4.
  • סעיף 3 - חשבון: block של 256 threads = 256/32 = 8 warps = 8/4 = 2 warpgroups. warpgroup 0 = warp-ranks 0,1,2,3; warpgroup 1 = warp-ranks 4,5,6,7.
  • סעיף 4 - שלושת היתרונות: (1) ריווי Tensor Cores - tile כפל-מטריצות בגודל 128 threads מזין את כל ארבעת ה-Tensor Cores עם רוחב פס אריתמטי גדול בהוראה אחת. (2) ביטול סנכרון בין-warp - ארבעת ה-warps משתפים פעולה מרומז ברמת ה-warpgroup, בלי __syncthreads ידני באמצע כפל המטריצה. (3) אסינכרוניות - wgmma.mma_async רץ ברקע, כך שה-warpgroup טוען את הtile הבא מ-global/shared בזמן שכפל המטריצה הנוכחי מתבצע. חפיפה זו היא בדיוק מה שמאפשר ל-FlashAttention על Hopper לחפוף את חישוב ה-attention (על ה-Tensor Cores) עם תנועת הנתונים, ולהגיע לניצול גבוה של החומרה.

למה זה עבד: המבנה של ה-warpgroup נגזר ישירות ממבנה ה-SM (ארבע תת-מחיצות), ולכן היחידה הזו "יושבת בול" על החומרה. איך להכליל: תמיד כשרואים יחידת תיאום חדשה (warp = 32, warpgroup = 128), שאלו לאיזו יחידת חומרה היא ממופה - זה מסביר מיד את הגודל ואת מגבלות היישור. את ה-kernels ברמת ה-Tensor Core וה-wgmma נכתוב בפרקים המתקדמים של הקורס.