לדלג לתוכן

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

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

בתרגול הזה נבנה reduction (צמצום) של מערך במספר גרסאות, ונרגיש בידיים את ההבדל בין תיאום ברמת block לתיאום ברמת warp. נתחיל בגרסה הקלאסית מבוססת shared memory + __syncthreads, נעבור לגרסת warp shuffles, ננקה אותה עם Cooperative Groups, נסגור עם atomic חוצה-blocks, ולבסוף ננתח למה מחסום מפוצל גורם ל-deadlock ולמה אין מחסום חוצה-blocks. עבדו לפי הסדר - כל גרסה בונה על הקודמת. אפשר לעבוד על כרטיס NVIDIA מקומי או ב-Google Colab (T4 חינם). אם אתם ב-Colab, החליפו -arch=sm_90 ב--arch=sm_75 והוסיפו ! לפני כל פקודת shell. שימו לב שכמה מהמספרים (למשל תקרת 64 warps ל-SM, ו-__reduce_*_sync שדורש CC 8.0+) הם של Ampere/Hopper; ל-Turing (T4, CC 7.5) אין __reduce_*_sync חומרתי, אבל כל שאר התרגילים עובדים.

הכנה - שלד ה-reduction

צרו את reduce.cu. הוא כולל את מאקרו בדיקת השגיאות, יצירת קלט אקראי, חישוב הסכום הנכון על ה-CPU להשוואה, ומקום לארבעה kernels שנמלא לאורך התרגילים.

// reduce.cu
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <cuda_runtime.h>
#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

#define CUDA_CHECK(call)                                                   \
    do {                                                                   \
        cudaError_t err_ = (call);                                         \
        if (err_ != cudaSuccess) {                                         \
            fprintf(stderr, "CUDA error at %s:%d: %s\n",                   \
                    __FILE__, __LINE__, cudaGetErrorString(err_));         \
            exit(EXIT_FAILURE);                                            \
        }                                                                  \
    } while (0)

__inline__ __device__ float warpReduceSum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;
}

// --- the kernels from the exercises go here ---

int main() {
    const int n = 1 << 24;                 // ~16.7 million elements
    const int block = 256;
    const int grid  = (n + block - 1) / block;
    size_t bytes = (size_t)n * sizeof(float);

    float* h_in = (float*)malloc(bytes);
    double ref = 0.0;
    for (int i = 0; i < n; ++i) { h_in[i] = 1.0f; ref += h_in[i]; }  // known sum = n

    float *d_in, *d_partial;
    CUDA_CHECK(cudaMalloc(&d_in, bytes));
    CUDA_CHECK(cudaMalloc(&d_partial, grid * sizeof(float)));
    CUDA_CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));

    // example launch (replace the kernel name in each exercise):
    reduceShared<<<grid, block, block * sizeof(float)>>>(d_in, d_partial, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    // combining the blocks' partial sums on the CPU (in exercise 4 we'll do this with atomic)
    float* h_partial = (float*)malloc(grid * sizeof(float));
    CUDA_CHECK(cudaMemcpy(h_partial, d_partial, grid * sizeof(float), cudaMemcpyDeviceToHost));
    double sum = 0.0;
    for (int i = 0; i < grid; ++i) sum += h_partial[i];

    printf("GPU sum = %.1f,  reference = %.1f,  %s\n",
           sum, ref, (fabs(sum - ref) < 1.0) ? "OK" : "MISMATCH");

    cudaFree(d_in); cudaFree(d_partial);
    free(h_in); free(h_partial);
    return 0;
}

הcompilation: nvcc -arch=sm_90 reduce.cu -o reduce והרצה ./reduce.

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

נממש את הגרסה הבסיסית.

  1. הוסיפו kernel בשם reduceShared שמבצע: טעינת global -> shared, __syncthreads, ואז לולאת צמצום עץ שבה בכל שלב s (החל מ-blockDim.x/2, מתחלק ב-2) חצי מה-threads מסכמים זוג, עם __syncthreads אחרי כל שלב. ה-thread מספר 0 כותב את sdata[0] אל out[blockIdx.x].
  2. ודאו ש-__syncthreads נמצא מחוץ ל-if (tid < s) - כל thread חייב להגיע אליו.
  3. קמפלו והריצו. ודאו שהפלט OK (הסכום שווה ל-n).
  4. ספרו: לכמה __syncthreads מגיעים ב-block של 256 threads (אחד אחרי הטעינה + אחד לכל שלב צמצום)?

רמז: לולאת הצמצום היא for (int s = blockDim.x/2; s > 0; s >>= 1), והגוף הוא if (tid < s) sdata[tid] += sdata[tid + s]; ואחריו __syncthreads();. גודל ה-shared memory הדינמי מועבר כארגומנט השלישי ב-<<<grid, block, block*sizeof(float)>>>.

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

עכשיו נחליף את רוב ה-shared memory וה-מחסומים בפרימיטיבי warp.

  1. הוסיפו kernel בשם reduceShfl שמשתמש ב-warpReduceSum (כבר בשלד) כדי לצמצם קודם כל warp, ואז כותב את סכום כל warp אל __shared__ float warpSums[32], מסנכרן פעם אחת עם __syncthreads, וה-warp הראשון מצמצם את הסכומים החלקיים.
  2. חשבו lane = threadIdx.x % warpSize ו-wid = threadIdx.x / warpSize. רק lane == 0 כותב ל-warpSums[wid].
  3. הריצו וודאו OK.
  4. ספרו כמה __syncthreads יש עכשיו ב-kernel. השוו למספר מתרגיל 1. כמה תאי shared memory משמשים כאן לעומת שם?

רמז: מספר ה-warps ל-block הוא blockDim.x / warpSize (עבור 256 -> 8). בצעד השני, רק wid == 0 פעיל, וטוענים val = (threadIdx.x < numWarps) ? warpSums[lane] : 0.0f לפני קריאה שנייה ל-warpReduceSum.

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

נמדוד את ההבדל בין שתי הגרסאות.

  1. עטפו כל launch של kernel ב-cudaEvent (cudaEventRecord לפני ואחרי, cudaEventElapsedTime) והדפיסו זמן ריצה במילישניות. הריצו כל kernel 100 פעמים בלולאה וקחו ממוצע (אחרי warmup של launch אחת).
  2. הריצו את שתי הגרסאות על אותו קלט. ודאו ששתיהן נותנות OK (אותה נכונות).
  3. רשמו את הזמנים. איזו גרסה מהירה יותר על ה-GPU שלכם, ובכמה אחוזים?
  4. הסבירו את ההפרש במונחי הסעיפים מההרצאה: כמה __syncthreads, כמה גישות shared, וכמה bank conflicts פוטנציאליים בכל גרסה.

רמז: לתזמון מדויק, סנכרנו עם cudaEventSynchronize על ה-event השני לפני cudaEventElapsedTime. ההפרש נובע מכך שגרסת ה-shuffle מתאמת באוגרים (זול), ולולאת ה-shared מצמצמת דרך ה-L1 עם מחסום בכל שלב (יקר). על קלט גדול ההפרש בולט יותר.

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

ננקה את הקוד ונסגור את הצעד החוצה-blocks בלי מעבר דרך ה-CPU.

  1. הוסיפו kernel reduceCG שמשתמש ב-cg::this_thread_block(), ב-cg::tiled_partition<32>(block), וב-cg::reduce(warp, val, cg::plus<float>()) במקום warpReduceSum הידני. השתמשו ב-warp.meta_group_rank() לאינדקס ה-warp וב-block.sync() במקום __syncthreads.
  2. במקום שכל block יכתוב לתא out[blockIdx.x] ואז נסכם על ה-CPU, תנו ל-thread מספר 0 של כל block לבצע atomicAdd(out, blockSum) אל תא גלובלי יחיד. אתחלו את התא ל-0 עם cudaMemset לפני הlaunch.
  3. הריצו וודאו שהתא הגלובלי היחיד מכיל את הסכום המלא (OK), בלי לולאת סיכום על ה-CPU.
  4. הסבירו: למה השתמשנו ב-atomicAdd ולא במחסום כדי לצרף את תוצאות ה-blocks? מה היה קורה אילו ניסינו לתאם את ה-blocks עם __syncthreads?

רמז: cg::reduce דורש את הכותרת <cooperative_groups/reduce.h>. ה-atomicAdd(float*, float) נתמך על כל כרטיס מ-CC 2.0 ומעלה. זכרו שאין דרך לגרום ל-block אחד לחכות ל-block אחר - ולכן atomic ל-global memory הוא הערוץ היחיד; הסדר בין ה-atomics אינו מובטח, אבל החיבור קומוטטיבי אז התוצאה נכונה.

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

נשתמש ב-Cooperative Groups כדי לכתוב צמצום נקי ברמת tile בגודל כלשהו.

  1. כתבו פונקציית __device__ גנרית template<int TILE> __device__ float tileReduce(cg::thread_block_tile<TILE> g, float v) שמבצעת return cg::reduce(g, v, cg::plus<float>());.
  2. ב-kernel, צרו tile בגודל 32 עם tiled_partition<32> וקראו לפונקציה. ואז נסו tile בגודל 16 (tiled_partition<16>) והשוו את התוצאה - כל tile של 16 מצטמצם בנפרד.
  3. הדפיסו מ-thread אחד לכל tile את g.thread_rank(), g.size(), ו-g.meta_group_rank() כדי לראות איך ה-block מתחלק.
  4. הסבירו: מה היתרון של tiled_partition על פני שימוש ישיר ב-__shfl_down_sync עם מסכה ידנית? מתי tile קטן מ-32 שימושי?

רמז: tiled_partition<N> דורש ש-N יהיה חזקה של 2 עד 32. הקבוצה נושאת את המסכה בתוכה, אז אין סכנה לטעות במסכה כשעובדים עם tiles חלקיים. ה-tiles הקטנים שימושיים כשכל קבוצה קטנה של threads מעבדת יחידת עבודה נפרדת (למשל 16 threads לכל שורה של מטריצה קטנה).

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

תרגיל ניתוח והדגמה. אזהרה: תרגיל זה יוצר deadlock מכוון - ה-kernel יתקע. הריצו עם timeout או היו מוכנים להרוג את התהליך.

  1. כתבו kernel badSync שבו __syncthreads() נמצא בתוך if (threadIdx.x < 100) { ... __syncthreads(); } על block של 256 threads.
  2. הריצו אותו (עם timeout 10 ./reduce או מוכנים ל-Ctrl+C). תארו מה קורה - האם ה-kernel חוזר? מה מחזיר cudaDeviceSynchronize?
  3. הסבירו במדויק, במונחי ספירת ה-threads שהגיעו למחסום: מדוע ה-block קופא. כמה threads הגיעו למחסום וכמה חסרים?
  4. תקנו את ה-kernel כך שכל ה-threads מגיעים ל-__syncthreads, וההבחנה בין ה-threads נעשית בערך ולא בזרימת הבקרה.

רמז: המחסום ממתין שכל 256 ה-threads יגיעו, אבל רק 100 נכנסו לענף. 156 ה-threads האחרים לעולם לא יגיעו, וה-100 שהגיעו ממתינים להם לנצח. התיקון: הוציאו את __syncthreads מחוץ ל-if, ותנו לתנאי לשלוט רק בערך (למשל float x = (threadIdx.x < 100) ? a : 0.0f;) ולא בהגעה למחסום.

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

תרגיל מחקר, בלי בהכרח קוד רץ (דורש H100/sm_90a).

  1. הריצו nvcc --list-gpu-arch וודאו אם sm_90a מופיע (התמיכה ב-wgmma דורשת את ה-a-suffix). אם אתם ב-Colab על T4, ציינו שהמושג לא רלוונטי שם ולמה.
  2. הסבירו במילים שלכם, מהמבנה של ה-SM (פרק 1): מדוע warpgroup הוא בדיוק 4 warps ולא 3 או 8, וכיצד הוא ממופה על ארבע תת-המחיצות של ה-SM.
  3. חשבו: block של 256 threads = כמה warps = כמה warpgroups? מהם ה-warp-ranks של כל warpgroup?
  4. הסבירו את שלושת היתרונות של wgmma.mma_async על פני כפל מטריצות מבוסס warp בודד: ריווי Tensor Cores, ביטול סנכרון בין-warp מפורש, ואסינכרוניות. קשרו ל-FlashAttention.

רמז: warpgroup ממפה warp אחד לכל תת-מחיצה, ולכל תת-מחיצה Tensor Core משלה - לכן 4. block של 256 = 8 warps = 2 warpgroups (warp-ranks 0-3 ו-4-7). ה-wgmma אסינכרוני, ולכן ה-warpgroup יכול לטעון את הtile הבא בזמן שכפל המטריצה הנוכחי רץ ברקע - חפיפה שהיא לב הביצועים של kernels כמו FlashAttention.