Cooperative Groups 深入 (Cooperative Groups Deep Dive)

重點總覽

4.4.1 介紹：為什麼需要 Cooperative Groups

Cooperative Groups 是 CUDA 程式模型的擴充，用來組織「協同合作的 thread 群組」，讓開發者控制 thread 協作的粒度（granularity），表達更豐富、更高效的平行分解，並內建常見平行原語（如 scan、parallel reduce）。

• 歷史上 CUDA 只提供單一簡單機制做 thread 同步：對整個 thread block 的 barrier，即 __syncthreads()。
• 為表達更廣的平行互動模式，許多人自行手刻warp 內或跨 block 的同步原語——效能雖有提升，卻累積出脆弱、難寫難調難維護、且跨 GPU 世代易壞的程式碼。
• Cooperative Groups 正是為此提供安全且 future-proof 的寫法。

4.4.2 Group Handle 與 Member Functions

Cooperative Group 透過 Cooperative Group Handle 管理。Handle 讓參與的 thread 得知自己在 group 中的位置、group 大小與其他資訊。

完整清單見 Cooperative Groups API。

4.4.3 預設行為 / Groupless Execution

代表 grid 與 thread block 的 group 會依 kernel launch 配置隱式建立。這些「implicit groups」是進一步細分的起點。

4.4.3.1 盡早建立 implicit group handle

為求最佳效能，建議在任何分支發生之前、盡可能早建立 implicit group 的 handle，並在整個 kernel 沿用同一個 handle。

4.4.3.2 只用 reference 傳遞 group handle

• 把 group handle 傳入函式時，建議以 reference 傳遞。
• Group handle 沒有 default constructor，必須在宣告時就初始化。
• 不鼓勵 copy-construct group handle。

namespace cg = cooperative_groups;
// 盡早取得 handle，並以 reference 傳給其他函式
cg::thread_block block = cg::this_thread_block();

重點：handle 是「輕量把手」但語意上不該到處複製，宣告即初始化、傳參用 reference。

4.4.4 建立 Cooperative Groups

Group 由切分（partition）母 group 而來；切分後會產生一個新 handle 管理子 group。

namespace cg = cooperative_groups;
// 取得目前 thread 的 cooperative group
cg::thread_block my_group = cg::this_thread_block();
// 切成大小為 8 的 tile
cg::thread_block_tile<8> my_subgroup = cg::tiled_partition<8>(my_group);
// 以 my_subgroup 進行運算

tiled_partition<8> 把 block 切成多個 8-thread 的 tile，每個 thread 取得管理自己 tile 的 handle。

  母 group (this_thread_block, 例: 32 threads)
  ┌──────────────────────────────────────────────────────────┐
  │ t0..t7 │ t8..t15 │ t16..t23 │ t24..t31                    │
  └──────────────────────────────────────────────────────────┘
   tiled_partition<8>  → 4 個固定大小子 group (1D, row-major)
   [tile0]    [tile1]     [tile2]     [tile3]

4.4.4.1 避免 Group 建立的 Hazards

4.4.5 同步（Synchronization）

在 Cooperative Groups 之前，CUDA 只允許 thread block 之間在 kernel 完成邊界才同步。Cooperative Groups 讓開發者可在不同粒度同步協作的 thread 群。

4.4.5.1 Sync

呼叫 collective 的 sync() 即可同步一個 group。與 __syncthreads() 相同，sync() 保證：

• 同步點之前group 內 thread 做的所有記憶體存取（讀/寫），對同步點之後的所有 thread 皆可見。
• group 內所有 thread 都到達同步點後，才允許任何 thread 越過。

namespace cg = cooperative_groups;
cg::thread_block my_group = cg::this_thread_block();
// 同步 block 內的 thread（等價 __syncthreads()）
cg::sync(my_group);

4.4.5.2 Barriers

Cooperative Groups 提供類似 cuda::barrier 的 barrier API 做更進階的同步，與 cuda::barrier 的關鍵差異：

• Cooperative Groups barrier 自動初始化。
• group 內所有 thread 每個 phase 都必須 arrive 並 wait 各一次。
• barrier_arrive 回傳一個 arrival_token，必須傳入對應的 barrier_wait，在那裡被消耗（consumed）且不可再用。

namespace cg = cooperative_groups;
cg::thread_block my_group = this_block();
cg::cluster_group cluster = this_cluster();
auto token = cluster.barrier_arrive();
// 選用：做些 local 處理以隱藏同步延遲
local_processing(my_group);
// 確保 cluster 內其他 block 都已執行並初始化 shared data 後，才存取 dsmem
cluster.barrier_waitmove(token);

barrier_arrive 與 barrier_wait 之間的空檔可拿來做不依賴同步結果的 local 工作，藉此隱藏同步延遲。

  phase 內每個 thread 的時序：
    barrier_arrive() ──► [token]
        │  (此空檔可做 local_processing，但不可做 collective op)
        ▼
    barrier_wait(move(token))   ← token 在此被消耗

4.4.6 集合運算（Collective Operations）

Collective 操作需 group 內所有 thread 參與才能完成。除非 API 明確允許，所有 thread 對對應參數必須傳入相同值，否則行為未定義。

4.4.6.1 Reduce

reduce 對 group 內每個 thread 提供的資料做平行歸約，須指定運算子：

namespace cg = cooperative_groups;
cg::thread_block my_group = cg::this_thread_block();
int val = data[threadIdx.x];
int sum = cg::reduce(my_group, val, cg::plus<int>());
if (my_group.thread_rank() == 0) {
    result[blockIdx.x] = sum;   // 由 rank 0 寫出結果
}

4.4.6.2 Scans

提供 inclusive_scan 與 exclusive_scan，適用任意 group 大小，對 group 內每個 thread 的資料做掃描；可選擇性指定上表的歸約運算子。

namespace cg = cooperative_groups;
cg::thread_block my_group = cg::this_thread_block();
int val = data[my_group.thread_rank()];
int exclusive_sum = cg::exclusive_scan(my_group, val, cg::plus<int>());
result[my_group.thread_rank()] = exclusive_sum;

exclusive_scan 不含自身、inclusive_scan 含自身，是 prefix-sum 類前綴運算的標準工具。

4.4.6.3 Invoke One

當「只需單一 thread 代表整個 group 做一段序列工作」時使用：

• invoke_one：從呼叫 group 中選一個任意 thread，用該 thread 以給定參數呼叫 invocable 函式。
• invoke_one_broadcast：同 invoke_one，但會把呼叫結果廣播給 group 內所有 thread。
• thread 選擇機制不保證 deterministic。

namespace cg = cooperative_groups;
cg::thread_block my_group = cg::this_thread_block();
// 確保 block 內只有一個 thread 印出訊息
cg::invoke_one(my_group, []() {
    printf("Hello from one thread in the block!");
});
// 同步讓所有 thread 等訊息印完
cg::sync(my_group);

4.4.7 非同步資料搬移（memcpy_async）

memcpy_async 提供 global ↔ shared memory 的非同步記憶體複製，特別適合最佳化記憶體傳輸、把計算與資料傳輸重疊以提升效能。

• memcpy_async：啟動從 global memory 到 shared memory 的非同步載入，當作「prefetch」使用——在資料被需要之前先載入。
• wait：強制 group 內所有 thread 等到非同步傳輸完成。所有 thread 都必須先呼叫 wait，才能存取 shared memory 中的資料。

namespace cg = cooperative_groups;
cg::thread_group my_group = cg::this_thread_block();
__shared__ int shared_data[];
// 從 global 非同步複製到 shared
cg::memcpy_async(my_group, shared_data + my_group.rank(),
                 input + my_group.rank(), sizeof(int));
// 此處做其他工作以隱藏延遲，但不可使用 shared_data
cg::wait(my_group);
// prefetch 的資料現在可用

  時間 ──►
  memcpy_async (global→shared) ████████████  (背景進行)
  其他計算 (不碰 shared_data)   ▓▓▓▓▓▓        ← 與傳輸重疊
  wait() ─ barrier ─ 之後才安全存取 shared_data

4.4.7.1 Memcpy Async 對齊需求

4.4.8 大規模群組（Large Scale Groups）

Cooperative Groups 支援橫跨整個 grid 的大型 group。前述所有功能對大型 group 皆可用，但有一個關鍵例外：

4.4.8.1 何時使用 cudaLaunchCooperativeKernel

• 用於啟動「需要 inter-block 同步」的單一 device kernel，使整個 grid 的所有 thread 都能跨 block 同步、協作。
• 保證 launch 是 atomic：若 API 呼叫成功，則指定數量的 thread block 必定會在指定 device 上啟動。
• 最佳實務：先查詢 device attribute cudaDevAttrCooperativeLaunch 確認支援。

int dev = 0;
int supportsCoopLaunch = 0;
cudaDeviceGetAttribute(&supportsCoopLaunch, cudaDevAttrCooperativeLaunch, dev);
// supportsCoopLaunch == 1 表示 device 0 支援 cooperative launch

支援條件：僅 Compute Capability 6.0 以上，且需在以下其中一種平台執行：

考試/測驗重點

Related Notes

• 04-CUDA-Features/11-Asynchronous-Barriers-Deep-Dive
• 04-CUDA-Features/12-Pipelines-Deep-Dive
• 04-CUDA-Features/13-Async-Copies-LDGSTS
• 04-CUDA-Features/08-Programmatic-Dependent-Launch-Deep-Dive
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 04-CUDA-Features/14-Async-Copies-TMA
• 04-CUDA-Features/15-Async-Copies-STAS
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps