2.4 סנכרון, cooperative groups ו warpgroups הרצאה

בשיעורים 2.1 ו-2.2 בנינו את היררכיית ה-threads וכתבנו kernel שבו כל thread עובד על נתח נתונים משלו, בבידוד מוחלט מהאחרים. זה מספיק לפעולות "מבוכות-מקבילית" (embarrassingly parallel) כמו חיבור וקטורים, אבל רוב האלגוריתמים המעניינים דורשים ש-threads ידברו זה עם זה: reduction (צמצום) של מיליון מספרים לסכום אחד, prefix-sum, היסטוגרמה, כפל מטריצות מבוזר-shared. ברגע ש-threads חולקים נתונים, נולדת בעיית תזמון קלאסית - מי כתב, מי קרא, ובאיזה סדר. השיעור הזה עוסק בכלי התיאום של CUDA לאורך כל ההיררכיה: המחסום __syncthreads() ברמת ה-block, פרימיטיבי ה-warp שמאפשרים ל-32 threads להחליף אוגרים בלי לגעת בזיכרון בכלל, ה-API המודרני של Cooperative Groups שנותן לכל אלה שם וטיפוס, ולבסוף ה-warpgroup של Hopper - יחידת התיאום בת 128 הthreads שמניעה את ה-kernels המהירים בעולם כמו FlashAttention. נראה שכל רמה בהיררכיה מקבלת ערוץ תקשורת משלה, ושהחוקים הנוקשים של הסנכרון אינם גחמה אלא נובעים ישירות ממה שראינו על החומרה בפרק 1.

המחסום ברמת ה-block - __syncthreads¶

הכלי הבסיסי ביותר לתיאום הוא מחסום (barrier) ברמת ה-block. הקריאה __syncthreads() היא נקודת מפגש: כל thread שמגיע אליה עוצר וממתין, עד שכל שאר ה-threads באותו block הגיעו גם הם. רק אז כולם משתחררים וממשיכים יחד. המחסום עושה שני דברים בבת אחת: הוא מסנכרן את ההתקדמות (כולם באותה נקודה בזמן), והוא משמש memory fence ברמת ה-block - כל כתיבה ל-shared memory או ל-global memory שבוצעה לפני המחסום מובטחת להיות גלויה לכל ה-threads של ה-block אחרי המחסום.

מדוע המחסום קיים רק בתוך block, ולא מעל? התשובה מגיעה ישירות מפרק 1. ראינו בשיעור 1.2 ובגלוסרי שכל ה-threads של block יושבים יחד על SM יחיד (block הוא ה-CTA - Cooperative Thread Array - מערך תהליכונים משתף-פעולה). המחסום הוא ספירה חומרתית זולה בתוך ה-SM: המתזמנים סופרים כמה warps של ה-block הגיעו, וכשכולם שם - משחררים. מכיוון שכל ה-block על אותו SM, הספירה מקומית ומהירה. block אחר יושב על SM אחר, אולי אפילו לא התחיל לרוץ עדיין (בשיעור 2.3 ראינו ש-blocks מתוזמנים בכל סדר), ולכן אין שום מנגנון חומרתי שיכול לחצות SMs ולסנכרן ביניהם. אין __syncthreads גלובלי כי אין חומרה שתממש אותו בזול.

הדפוס הקנוני שבו משתמשים ב-__syncthreads הוא global -> shared -> syncthreads -> compute: כל ה-block טוען יחד נתח נתונים מ-global memory אל shared memory, מסנכרן כדי לוודא שהנתח נטען במלואו, ורק אז מעבד. זהו בדיוק דפוס ה-tiling שנעמיק בו בפרק 3.3.

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

    tile[tid] = (i < n) ? in[i] : 0.0f;   // 1. global -> shared (each thread writes one cell)
    __syncthreads();                       // 2. barrier: the whole tile is loaded

    // 3. compute: now safe to read neighbor cells written by another thread
    float left  = (tid > 0)   ? tile[tid - 1] : 0.0f;
    float right = (tid < 255) ? tile[tid + 1] : 0.0f;
    if (i < n) out[i] = 0.25f*left + 0.5f*tile[tid] + 0.25f*right;
}

בלי המחסום בשורה 2, thread מספר 5 עלול לקרוא את tile[6] לפני ש-thread מספר 6 הספיק לכתוב אליו - תוצאה שגויה ולא דטרמיניסטית (race condition, מרוץ נתונים). המחסום הופך את הכתיבות של כולם לגלויות לכולם.

חוק הזהב - אסור מחסום מפוצל¶

יש כלל נוקשה אחד סביב __syncthreads, וזהו מקור הבאגים הנפוץ ביותר בקורס הזה: כל ה-threads של ה-block חייבים להגיע לאותו __syncthreads בדיוק. אסור שחלק מה-threads יגיעו למחסום וחלק יפסחו עליו. אם הם עושים זאת, ה-block נכנס ל-deadlock (קיפאון) - ה-threads שהגיעו למחסום ממתינים לנצח ל-threads שלעולם לא יגיעו, וה-kernel תקוע.

הדוגמה הקלאסית לשגיאה היא מחסום בתוך ענף מפוצל (divergent branch):

// wrong! deadlock if n is not a multiple of blockDim.x
if (i < n) {
    tile[tid] = in[i];
    __syncthreads();   // only some threads enter the branch and reach here
}

// correct: every thread reaches the barrier, and the bounds guard happens in the load itself
tile[tid] = (i < n) ? in[i] : 0.0f;
__syncthreads();

בגרסה השגויה, אם n אינו מתחלק בגודל ה-block, ה-threads האחרונים (שעבורם i >= n) לא נכנסים לענף if, לא מגיעים למחסום, וה-block קופא. הכלל הפרקטי: קריאות __syncthreads חייבות להיות מחוץ לכל תלות-נתונים בזהות ה-thread. הגנות גבול מבצעים בערך שנטען (טוענים 0.0f מחוץ לתחום), לא בדילוג על המחסום.

הערה חשובה על גרנולריות ה-warp: מכיוון שהחומרה מתזמנת warps שלמים, __syncthreads שמפוצל ברמת ה-warp (warp שלם נכנס או לא נכנס) עלול לפעמים "לעבוד" במקרה על ארכיטקטורות ישנות, אבל זו התנהגות בלתי-מוגדרת (undefined behavior) שמתפרקת עם Independent Thread Scheduling של Volta ומעלה. אל תסתמכו עליה לעולם.

ערוצי התקשורת בהיררכיה - shared, global ו-atomics¶

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

בתוך warp: ה-threads יכולים להחליף ערכי אוגרים ישירות דרך פרימיטיבי shuffle (הסעיף הבא), בלי לגעת בזיכרון כלל. זה הערוץ המהיר ביותר.
בתוך block: הערוץ הוא shared memory (זיכרון משותף), שראינו בשיעור 1.6 שהוא פיזית ה-L1 data cache של ה-SM (ב-H100, עד 228 KiB shared מתוך 256 KiB L1 ל-SM). התיאום נעשה עם __syncthreads. זה מהיר - עשרות מחזורים - כי הכל על השבב ועל אותו SM.
חוצה blocks: הערוץ היחיד הוא global memory בשילוב פעולות atomic (אטומיות). אין מחסום חוצה-blocks (ראינו למה), ולכן blocks שונים מתאמים רק דרך קריאה/כתיבה אטומית ל-global memory, למשל atomicAdd. זה איטי - מאות מחזורים - וללא סדר מובטח.

  hierarchy level    communication channel    sync tool            cost
  ------------        -----------              -----------          ----
  warp (32)            registers via shuffle    __syncwarp           ~a few cycles
  block (CTA)          shared memory (L1)       __syncthreads        tens of cycles
  grid                 global memory            atomics (no barrier) hundreds of cycles

הטבלה הזו היא מפת הדרכים לתכן אלגוריתמים: תמיד שואפים לדחוף כמה שיותר תיאום למטה בהיררכיה (warp לפני block, block לפני grid), כי כל צעד למטה זול בסדר גודל. reduction טובה, למשל, מצמצמת קודם בתוך כל warp, אז בתוך ה-block, ורק בסוף - צעד אחד קטן - חוצה blocks דרך atomic.

פרימיטיבי ה-warp - warp-level primitives¶

בתוך warp החומרה כבר מריצה 32 threads ב-lockstep (צעד-נעול) על אותה הוראה. זה פותח ערוץ תקשורת ייחודי: ה-threads יכולים לקרוא ישירות זה את אוגרי זה, בלי shared memory ובלי מחסום. משפחת ההוראות שעושה זאת נקראת shuffle, וכל הפרימיטיבים האלה נושאים סיומת _sync וארגומנט mask (מסכה) שמציין אילו lanes (נתיבים - מיקום ה-thread בתוך ה-warp, 0..31) משתתפים. המסכה הנפוצה 0xffffffff פירושה "כל 32 ה-threads". ה-_sync מבטיח שכל ה-threads שבמסכה מגיעים להוראה יחד לפני שהחילוף מתבצע - זהו מחסום זעיר ברמת ה-warp.

פרימיטיב	מה עושה
`__syncwarp(mask)`	מחסום ברמת warp: מסנכרן את ה-lanes שבמסכה בלבד
`__shfl_sync(mask, v, src)`	כל lane מקבל את `v` של lane מספר `src` (שידור/החלפה כללית)
`__shfl_down_sync(mask, v, d)`	כל lane מקבל את `v` של ה-lane שנמצא `d` מקומות מעליו
`__shfl_up_sync(mask, v, d)`	כל lane מקבל את `v` של ה-lane שנמצא `d` מקומות מתחתיו
`__shfl_xor_sync(mask, v, m)`	כל lane מחליף עם ה-lane שמספרו `laneid XOR m` (butterfly)
`__ballot_sync(mask, p)`	מחזיר מסכת 32-bit: ביט i דולק אם ה-predicate `p` של lane i אמת
`__any_sync(mask, p)`	מחזיר אמת אם `p` אמת אצל לפחות lane אחד
`__all_sync(mask, p)`	מחזיר אמת אם `p` אמת אצל כל ה-lanes
`__reduce_add_sync(mask, v)`	צמצום חומרתי (סכום) על פני ה-warp - רק CC 8.0+, ורק שלמים

הפרימיטיב החשוב ל-reduction הוא __shfl_down_sync. הוא מאפשר צמצום עץ (tree reduction) בתוך warp בלוג(32)=5 צעדים, בלי shared memory ובלי __syncthreads:

__inline__ __device__ float warpReduceSum(float val) {
    // reduction tree: 16, 8, 4, 2, 1 - five steps for a full warp
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;   // the full result sits in lane 0
}

  warp reduction with __shfl_down_sync (shown for 8 lanes; a real warp of 32 -> 5 steps)
  lane:   0    1    2    3    4    5    6    7
  start:  a0   a1   a2   a3   a4   a5   a6   a7
  off=4:  a0+a4 a1+a5 a2+a6 a3+a7  .    .    .    .
  off=2:  +a2.. +a3..   .    .
  off=1:  SUM    .
          ^ the full sum of all eight sits in lane 0

היתרון עצום: אין גישה ל-shared memory (חוסך bank conflicts - התנגשויות בנקים, ראו 3.3), אין __syncthreads, והכל בתוך אוגרים. __ballot_sync ו-__any/__all_sync שימושיים לקבלת החלטות קולקטיביות - למשל "האם למישהו ב-warp יש עבודה נוספת?" - ולדחיסת נתונים (stream compaction). על CC 8.0 ומעלה (H100 הוא 9.0), __reduce_add_sync עושה את הצמצום השלם בהוראה חומרתית אחת, אך רק לטיפוסים שלמים.

דוגמה מלאה - block reduction בשתי דרכים¶

נחבר הכל לבעיה אמיתית: סכימת מערך. נשווה שתי מימושי block 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();

    // tree reduction in shared: at each step half the threads sum a pair
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();                    // barrier after every step - mandatory
    }
    if (tid == 0) out[blockIdx.x] = sdata[0];
}

שימו לב ש-__syncthreads אינו בתוך ה-if (tid < s) אלא מחוצה לו - כל ה-threads מגיעים אליו בכל איטרציה, גם אלה ש-tid >= s. זה בדיוק חוק הזהב מהסעיף הקודם. ל-block של 256 threads יש 8 איטרציות (8 מחסומים).

הגרסה השנייה משתמשת ב-warp shuffles ומצמצמת קודם בכל warp, ואז מצרפת את הסכומים החלקיים:

__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);               // step 1: each warp reduces, result in lane 0

    __shared__ float warpSums[32];          // at most 1024/32 = 32 warps per block
    int lane = threadIdx.x % warpSize;
    int wid  = threadIdx.x / warpSize;
    if (lane == 0) warpSums[wid] = val;     // each warp writes one cell
    __syncthreads();                        // the only barrier in the whole kernel

    // step 2: the first warp reduces the partial sums
    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;
}

השוו את השתיים. נכונות: שתיהן נכונות, אבל הגרסה השנייה משתמשת ב-__syncthreads בדיוק פעם אחת במקום 8 פעמים, וכמעט לא נוגעת ב-shared memory (רק 32 תאים במקום כל ה-block). ביצועים: הצמצום התוך-warp רץ באוגרים בלבד, ללא bank conflicts וללא סנכרון-block, ולכן על H100 גרסת ה-shuffle מהירה משמעותית ל-block גדול. זו הסיבה שכל ספריות ה-reduction המודרניות (למשל CUB) בנויות סביב warp shuffles, ומורידות את התיאום כמה שיותר למטה בהיררכיה. את הצעד האחרון, צירוף התוצאות של כל ה-blocks, נבצע עם atomicAdd ל-global memory (אין מחסום חוצה-blocks), ונרחיב על כל הטכניקה בשיעור 9.4.

ה-API המסודר - Cooperative Groups¶

הפרימיטיבים שראינו הם פונקציות מובנות עם ארגומנט mask "עירום" - קל לטעות במסכה, וקשה לכתוב קוד גנרי שעובד גם על tile של 16 threads וגם על warp שלם. ספריית Cooperative Groups (קבוצות משתפות-פעולה), שנכללת ב-<cooperative_groups.h>, נותנת לכל רמת תיאום טיפוס מפורש עם מתודות .sync(), .thread_rank() ו-.size(). הקבוצה נושאת את המסכה בתוכה, אז אי אפשר לטעות בה.

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

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

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

    val = cg::reduce(warp, val, cg::plus<float>());            // clean warp reduction, no masks

    __shared__ float warpSums[32];
    if (warp.thread_rank() == 0)
        warpSums[warp.meta_group_rank()] = val;               // meta_group_rank = the warp's index within the block
    block.sync();                                              // equivalent to __syncthreads

    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) out[blockIdx.x] = val;
    }
}

הקוד הזה זהה בתוצאה ל-reduceShfl, אבל קריא בהרבה: block.sync() במקום __syncthreads(), ו-cg::reduce(warp, val, cg::plus<float>()) במקום לולאת ה-shuffle הידנית (על CC 8.0+ הוא מתקמפל ל-__reduce-חומרתי כשאפשר). הפונקציות המרכזיות:

cg::this_thread_block() - הקבוצה של ה-block הנוכחי (ה-CTA).
cg::tiled_partition<N>(g) - מחלקת קבוצה ל-tiles בגודל N בזמן קומפילציה (32 = warp, אבל גם 16, 8, 4...); כל tile הוא קבוצה עצמאית עם .sync() משלו.
cg::coalesced_threads() - מחזירה קבוצה של ה-threads הפעילים בפועל ברגע זה (אלה שלא פוצלו החוצה בענף). שימושי בתוך קוד מפוצל, כשלא ידוע מראש מי פעיל.
cg::this_grid() - הקבוצה של ה-grid כולו (ראו הסעיף הבא).

סנכרון חוצה-grid - grid_group והlaunch הקואופרטיבית¶

אמרנו שאין מחסום חוצה-blocks. זה נכון עבור השקת kernel רגילה. Cooperative Groups מציעה דלת אחורית מוגבלת: cg::this_grid() מחזירה grid_group, וקריאה ל-grid.sync() מסנכרנת את כל ה-blocks ב-grid. אבל יש מלכוד עמוק שנובע ישירות ממה שראינו על תזמון blocks.

כדי ש-grid.sync() לא ייכנס ל-deadlock, כל ה-blocks של ה-grid חייבים להיות תושבים (resident) על ה-SMs בו-זמנית. אם יש יותר blocks ממה שנכנס פיזית על ה-GPU, blocks מסוימים ממתינים בתור לתזמון - וה-blocks שכבר רצים תקועים ב-grid.sync() וממתינים להם, אבל הם לעולם לא ירוצו כי אין SM פנוי. deadlock. לכן launch קואופרטיבית דורשת נתיב מיוחד, cudaLaunchCooperativeKernel, ומחייבת לחשב מראש grid שנכנס כולו:

// mandatory: compute how many blocks fit simultaneously, and don't exceed it
int numBlocksPerSm = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocksPerSm, myKernel, blockSize, /*dynamicSMem=*/0));
int numSm = 0;
CUDA_CHECK(cudaDeviceGetAttribute(&numSm, cudaDevAttrMultiProcessorCount, 0));
dim3 grid(numBlocksPerSm * numSm);      // the maximum grid that fits at once

void* args[] = { &d_in, &d_out, &n };
CUDA_CHECK(cudaLaunchCooperativeKernel((void*)myKernel, grid, blockSize, args, 0, 0));

ה-CUDA_CHECK הוא מאקרו בדיקת שגיאות שנשתמש בו לאורך הקורס - הוא עוטף כל קריאת CUDA, בודק את קוד השגיאה המוחזר, ומדפיס קובץ+שורה אם משהו נכשל. במקרה של launch קואופרטיבית הוא קריטי, כי cudaErrorCooperativeLaunchTooLarge היא בדיוק השגיאה שנקבל אם ביקשנו grid גדול מדי. הכלל: grid.sync() הוא כלי נדיר ויקר; לרוב עדיף לפצל ל-שני kernels נפרדים (סנכרון בין launches הוא מחסום גלובלי "חינם"), ולהשתמש ב-grid group רק כשהאלגוריתם באמת דורש מספר סבבים על אותם נתונים תושבים.

היחידה של Hopper - warpgroup¶

הרמה הגבוהה ביותר בהיררכיית התיאום החומרתית הופיעה עם ארכיטקטורת Hopper (דור ה-H100): ה-warpgroup. warpgroup הוא אוסף של ארבעה warps רצופים, כך שה-warp-rank (הדירוג) של הראשון הוא כפולה של 4. זהו 4 x 32 = 128 threads. הדירוג מחושב מהאינדקס הליניארי של ה-thread ב-block:

  linearIdx = tid.x + tid.y*ntid.x + tid.z*ntid.x*ntid.y
  warpRank  = linearIdx / 32
  for a dispatch of 8 warps:
    warpgroup 0 = warp-ranks 0,1,2,3
    warpgroup 1 = warp-ranks 4,5,6,7

מדוע דווקא ארבעה, ודווקא ביישור לכפולה של 4? זה מתחבר ישירות למבנה ה-SM מפרק 1: ראינו ש-SM של H100 בנוי מארבע תת-מחיצות (sub-partitions), כל אחת עם warp scheduler משלה ו-Tensor Core (ליבת טנזור) משלה. warpgroup ממופה בול על ארבע התת-מחיצות - warp אחד לכל תת-מחיצה - כך שהוראה ברמת ה-warpgroup מפעילה את כל ארבעת ה-Tensor Cores של ה-SM בו-זמנית. (הגלוסרי מציין שהיישור-לכפולה-של-4 אינו מתועד רשמית, וזו השערה מושכלת המבוססת על מבנה ארבע התת-מחיצות.)

הסיבה שכל זה קיים היא הוראה אחת: wgmma.mma_async (warpgroup matrix multiply-accumulate, asynchronous - כפל-וצבירת-מטריצות אסינכרוני ברמת warpgroup). במקום ש-32 threads של warp בודד ישתפו פעולה על tile מטריצה קטן, 128 threads של warpgroup שלם מבצעים tile כפל-מטריצות גדול בהוראה אחת. שלושה יתרונות נובעים מכך:

ריווי ה-Tensor Cores: ה-tiles הגדולים יותר מזינים את רוחב הפס האריתמטי האדיר של ה-Tensor Cores טוב יותר - הוראה אחת שמעסיקה את כל ארבעת ה-Tensor Cores של ה-SM.
ביטול סנכרון בין-warp מפורש: ארבעת ה-warps משתפים פעולה מרומז ברמת ה-warpgroup, ולכן המתכנת אינו צריך __syncthreads ידני כדי לתאם ביניהם באמצע כפל המטריצה. התיאום מובנה בהוראה.
אסינכרוניות: ה-wgmma הן אסינכרוניות - כפל המטריצה מתקדם ברקע בעוד ה-warpgroup מבצע עבודה אחרת (למשל טוען את הtile הבא), מה שמאפשר חפיפה מלאה בין חישוב לזיכרון.

זו הסיבה שה-warpgroup קריטי ל-kernels המהירים ביותר של ימינו - כמו FlashAttention על Hopper. חשוב לזכור: warpgroup הוא מושג ספציפי ל-Hopper ומעלה. לפני Hopper הוא אינו יחידת תזמון ממדרגה ראשונה, וקוד שמסתמך על wgmma דורש compute capability 9.0 (ה-a-suffix, sm_90a, שראינו בשיעור 2.5). ברוב הקורס נתאם ברמת block ו-warp; את ה-warpgroup נפגוש שוב כשנגיע ל-kernels ברמת ה-Tensor Core בפרקים המתקדמים.

סיכום¶

הכלי הבסיסי לתיאום הוא __syncthreads() - מחסום ברמת ה-block שגם מסנכרן התקדמות וגם משמש memory fence; הוא קיים רק בתוך block כי כל ה-block יושב על SM יחיד, ואין חומרה שתסנכרן זול בין SMs.
הדפוס הקנוני הוא global -> shared -> __syncthreads -> compute: טוענים tile מ-global אל shared, מסנכרנים, ורק אז מעבדים - אחרת נוצר race condition.
חוק הזהב: כל ה-threads של ה-block חייבים להגיע לאותו __syncthreads; מחסום בתוך ענף מפוצל (divergent) גורם ל-deadlock, ולכן הגנות גבול עושים בערך הנטען ולא בדילוג על המחסום.
לכל רמה בהיררכיה ערוץ תקשורת יחיד: warp מחליף אוגרים דרך shuffle, block מתקשר דרך shared memory (פיזית ה-L1), וחוצה-blocks אפשרי רק דרך global memory + atomics - כל צעד למטה זול בסדר גודל.
פרימיטיבי ה-warp (__shfl_*_sync, __ballot_sync, __any/__all_sync, ועל CC 8.0+ גם __reduce_*_sync) מאפשרים ל-32 threads לתאם דרך אוגרים בלבד, בלי shared memory ובלי מחסום; __shfl_down_sync נותן צמצום warp ב-5 צעדים.
ה-reduction מבוססת shuffle משתמשת ב-__syncthreads פעם אחת בלבד וכמעט לא נוגעת ב-shared memory, ולכן מהירה מהגרסה הקלאסית מבוססת-shared; זו הסיבה שספריות כמו CUB בנויות סביב warp shuffles.
ספריית Cooperative Groups נותנת לכל רמת תיאום טיפוס מפורש (this_thread_block, tiled_partition<32>, coalesced_threads) שנושא את המסכה בתוכו, עם .sync() ו-cg::reduce שמחליפים את הפרימיטיבים העירומים בקוד קריא וגנרי.
סנכרון חוצה-grid קיים דרך grid.sync() על this_grid(), אבל דורש cudaLaunchCooperativeKernel ומחייב שכל ה-blocks יהיו תושבים בו-זמנית, אחרת deadlock - לכן לרוב עדיף לפצל לשני kernels.
ה-warpgroup של Hopper הוא ארבעה warps רצופים (128 threads) הממופים על ארבע תת-המחיצות של ה-SM; הוא מפעיל את wgmma.mma_async - כפל מטריצות אסינכרוני גדול שמרווה את כל ה-Tensor Cores, מבטל סנכרון בין-warp מפורש, ומניע kernels כמו FlashAttention.