Enter your search text in the box above

Select a result to preview

記憶體屏障、Pipeline 與 Cooperative Groups API (Memory Barrier, Pipeline and Cooperative Groups APIs)

#cuda #language-extensions #async-execution #cooperative-groups #async-barrier #concept

重點總覽

本篇整理三組 device-callable 底層 API：<cuda_awbarrier_primitives.h> 的 memory barrier C-like primitives（__mbarrier_*）、<cuda_pipeline.h> 的 pipeline primitives（__pipeline_*），以及 <cooperative_groups.h> 的 Cooperative Groups（群組型別 + collective 運算）。前兩者是 cuda::barrier/cuda::pipeline 的精簡 C 介面；CG 則用群組物件統一表達 thread block / cluster / grid / tile / coalesced，並提供 sync/shfl/reduce/scan 等 collective。

項目 重點
__mbarrier_* cuda::barrier 的 C-like primitives；init/arrive/test_wait/try_wait；bar 必須在 __shared__
__mbarrier_arrive vs arrive_and_drop 後者額外遞減「下一 phase」的 expected count
__pipeline_memcpy_async global→shared 非同步複製；size_and_align 只能是 4/8/16
__pipeline_commit / wait_prior(N) 提交批次 / 等到批次 L-N（含）完成
thread_block / cluster_group / grid_group 三大 implicit group；分別由 this_thread_block()/this_cluster()/this_grid() 取得
tiled_partition<Size> 切成 1D row-major 子群；Size 為 2 的次方且 ≤ 1024
coalesced_group coalesced_threads() 機會性取得當下 active threads
collective sync/memcpy_async/wait/reduce/inclusive_scan/barrier_arrive；全員必須參與，否則 UB
grid.sync() 需用 cudaLaunchCooperativeKernel，非 <<<>>>；CC 6.0+

Memory Barrier Primitives Interface

<cuda_awbarrier_primitives.h> 提供 cuda::barrier 功能的 C-like primitives。

資料型別

typedef /* implementation defined */ __mbarrier_t;       // barrier 物件
typedef /* implementation defined */ __mbarrier_token_t; // arrival token

API 一覽

函式 作用 / 條件
uint32_t __mbarrier_maximum_count() 回傳允許的最大 arrival count
void __mbarrier_init(__mbarrier_t* bar, uint32_t expected_count) bar 必須指向 __shared__；expected_count <= __mbarrier_maximum_count()；設定當前與下一 phase 的 expected arrival count
void __mbarrier_inval(__mbarrier_t* bar) 重新利用該 shared memory 前必須先 invalidate
__mbarrier_token_t __mbarrier_arrive(__mbarrier_t* bar) 原子遞減當前 phase 的 pending count；回傳遞減前狀態的 token
__mbarrier_token_t __mbarrier_arrive_and_drop(__mbarrier_t* bar) 原子遞減當前 phase 的 pending count 與下一 phase 的 expected count；回傳遞減前 token
bool __mbarrier_test_wait(__mbarrier_t* bar, __mbarrier_token_t token) 非阻塞測試；token 屬於「緊鄰前一 phase」則回 true，否則 false
bool __mbarrier_test_wait_parity(__mbarrier_t* bar, bool phase_parity) 以 phase 奇偶判定；true=奇數 phase，false=偶數 phase
bool __mbarrier_try_wait(__mbarrier_t* bar, __mbarrier_token_t token, uint32_t max_sleep_nanosec) 可能掛起執行緒；phase 完成回 true，超過時限回 false
bool __mbarrier_try_wait_parity(__mbarrier_t* bar, bool phase_parity, uint32_t max_sleep_nanosec) 同上，以 phase parity 判定
arrive 前提

呼叫 __mbarrier_arrive / __mbarrier_arrive_and_drop 前，*bar 必須已 init，且 pending count 不可為 0。

test_wait vs try_wait

test_wait 只是瞬間查詢（立即回傳）；try_wait 在尚未完成時可能讓執行緒睡眠，直到 phase 完成或達到 max_sleep_nanosec（取代系統預設時限）的奈秒時限。

Pipeline Primitives Interface

<cuda_pipeline.h> 提供 <cuda/pipeline> 的 C-like 介面；若編譯時無 ISO C++ 2011 相容，改 include <cuda_pipeline_primitives.h>。

適用範圍

Pipeline primitives 只追蹤 global → shared 的非同步複製，且有特定大小/對齊要求。它等同於一個 cuda::thread_scope_thread 的 cuda::pipeline 物件。

函式 作用
void __pipeline_memcpy_async(void* dst_shared, const void* src_global, size_t size_and_align, size_t zfill=0) 提交非同步複製：複製 size_and_align - zfill bytes，其餘 zfill bytes 補零
void __pipeline_commit() 把已提交的 memcpy_async 當作當前批次送進 pipeline
void __pipeline_wait_prior(size_t N) 設提交索引為 {0,1,...,L}，等待批次到 L-N（含）為止完成
void __pipeline_arrive_on(__mbarrier_t* bar) 對 barrier arrival count +1；此前所有 memcpy_async 完成時再 −1（淨效果為 0）

__pipeline_memcpy_async 語意（虛擬碼）：

size_t i = 0;
for (; i < size_and_align - zfill; ++i)            // 複製
    ((char*)dst_shared)[i] = ((char*)src_global)[i];
for (; i < size_and_align; ++i)                    // 補零
    ((char*)dst_shared)[i] = 0;

要求與限制：

race condition

從提交 memcpy_async 到 wait 完成之間，下列行為都會造成 race：

  • 讀取 dst_shared
  • 寫入 dst_shared 或 src_global
  • 對 dst_shared 或 src_global 做 atomic 更新

另外 __pipeline_arrive_on 的 +1 不可使 arrival count 超過 __mbarrier_maximum_count()，需使用者自行保證。

Cooperative Groups：群組型別

<cooperative_groups.h> 用群組物件統一表達各層級的執行緒集合。所有 collective 都要求群組內全員參與，否則行為未定義。

群組型別 建構方式 代表 條件
thread_block this_thread_block() 一個 thread block 衍生自更泛用的 thread_group
cluster_group this_cluster() 一個 cluster 的全部執行緒 CC 9.0+；非 cluster grid 視為 1×1×1 cluster
grid_group this_grid() 一個 grid 的全部執行緒 除 sync() 外 API 隨時可用；跨 grid 同步需 cooperative launch
thread_block_tile<Size, ParentT> tiled_partition<Size>(g) 編譯期固定大小的 tile Size 為 2 的次方且 ≤ 1024
coalesced_group coalesced_threads() 當下 active（coalesced）的執行緒 機會性、不保證成員

thread_block / cluster_group / grid_group 成員

共通：sync() 等同 barrier_wait(barrier_arrive())；barrier_arrive() 回傳 token 傳入 barrier_wait(arrival_token&&)；thread_rank()、num_threads()；legacy alias size() = num_threads()。thread_block 另有 legacy alias group_dim() = dim_threads()。

特有成員 群組 說明
group_index() / thread_index() / dim_threads() thread_block block 在 grid 的 3D index / thread 在 block 的 3D index / block 維度
block_rank() / num_blocks() / dim_threads() / dim_blocks() / block_index() cluster_group 在 cluster 內的 block 排名/總數；dim_threads()=cluster 的 thread 維度、dim_blocks()=block 維度；block_index()=block 在 cluster 的 3D index
query_shared_rank(addr) cluster_group 取得某 shared address 所屬的 block rank
map_shared_rank(addr, rank) cluster_group 取得 cluster 中另一 block 的 shared 變數位址
is_valid() grid_group 此 grid 是否可同步
block_rank() / cluster_rank() / num_blocks() / num_clusters() grid_group 回傳型別為 unsigned long long
dim_blocks() / dim_clusters() / block_index() / cluster_index() grid_group 回傳 dim3：grid 的 block/cluster 維度、block/cluster 在 grid 的 3D index；group_dim() = dim_blocks() 
__global__ void kernel(int *globalInput) {
    __shared__ int x;
    thread_block g = this_thread_block();
    if (g.thread_rank() == 0)   // 選一個 leader
        x = (*globalInput);     // 載入 shared 供全員使用
    g.sync();                   // 等同 __syncthreads();
}

thread_block_tile

函式 tile 尺寸限制
shfl 全尺寸；> 32 時所有 thread 必須給相同 src_rank，否則 UB
shfl_up / shfl_down / shfl_xor 僅 ≤ 32
ballot / match_any / match_all 僅 ≤ 32
非整數型別的 shuffle（C++11+）

只要 is_trivially_copyable<T>::value == true，且 sizeof(T) <= 32（tile ≤ 32）或 sizeof(T) <= 8（更大 tile），即可 shuffle 非整數型別。

大 tile 與舊硬體

Size > 32 在 CC 7.5 或更低需額外步驟：用 cooperative_groups::block_tile_memory<MaxBlockSize=1024>（置於 shared 或 global memory）並傳入 this_thread_block(shared)（此 overload 為 collective，須全 block 呼叫）。在 CC 8.0+ 可省略 block_tile_memory，以「同一份原始碼對應多種 CC」。

thread_block block = this_thread_block();
thread_block_tile<32> tile32 = tiled_partition<32>(block);       // provenance 存在 handle
thread_block_tile<4, thread_block> tile4 = tiled_partition<4>(block); // provenance 編進型別

coalesced_group

// 假設分支內只有 thread 2,4,8 為 active，coalesced_threads() 會建立
// 含 3 個 thread（rank 0-2）的群組
if (threadIdx.x == *globalInput) {
    coalesced_group active = coalesced_threads();
    active.sync();
}

Cooperative Groups：切分 (partition)

<cooperative_groups/partition.h>，三種 collective 切分：

API 行為 條件
tiled_partition<Size>(g) / tiled_partition(thread_group, tilesz) 切成 1D、row-major、大小相等的子群（共 size(parent)/tilesz 個） 父群大小須整除 Size；native 1/2/4/8/16/32；templated 另支援 64/128/256/512（CC ≤ 7.5 需額外步驟）。CC 5.0 起；> 32 需 C++11
labeled_partition(g, label) 依 label 值把相同 label 的 thread 分到同群 label 為整數型別。CC 7.0 起、C++11；功能仍在評估，未來可能微調
binary_partition(g, pred) labeled_partition 的特例，label 只能 0 或 1 labeled_partition 的特例（label 限 0/1）；CC 7.0／C++11 為比照 labeled 的延伸推論 
// 把 32-tile 切成奇數組與偶數組
auto tile32 = cg::tiled_partition<32>(block);
int elem = inputArr[block.thread_rank()];
auto subtile = cg::binary_partition(tile32, (elem & 1));

Cooperative Groups：Collective 運算

memcpy_async / wait（async.h）

Errata 與對齊

帶 dstLayout/srcLayout 的 overload 以元素而非 bytes 計數；元素型別由 TyElem 推得。若用 cuda::aligned_size_t<N>，對齊保證至少 min(16, N)，且 src/dst 都需對齊 N、複製 bytes 為 N 的倍數。Codegen：CC 5.0 起、CC 8.0 才有非同步性、C++11，需 include cooperative_groups/memcpy_async.h。

// 雙緩衝 pipeline：發出 stage N 的同時，等待並處理 stage N-1
cg::memcpy_async(tb, local_smem[stage ^ 1], elementsInShared,
                 global_data + index, elementsPerThreadBlock - index);
cg::wait_prior<1>(tb);   // 留最新 1 個還在跑，等其餘完成
// 處理 local_smem[stage] ...

reduce / scan（reduce.h、scan.h）

dispatch 行為 結果
cg::plus<int>() CC 8.0+ 特化為 __reduce_add_sync(...)（硬體）
cg::plus<float>() 無對應加速器，退回 shuffle-based reduction
vector / 自訂型別、lambda 無法被檢視以 dispatch，一律 shuffle-based
_update / _async 變體

reduce_update_async / reduce_store_async、inclusive_scan_update / exclusive_scan_update 接受 cuda::atomic 或 cuda::atomic_ref，把群組所有輸入之和按 op 以 relaxed ordering 原子更新到目標。atomic 的 scope 必須涵蓋群組所有 thread（多群共用同一 atomic 時須涵蓋全部）。要看到 async 結果，須同步呼叫群組或更大的包含群組。

  • reduce_*_async 為 void，只把結果原子更新到 atomic、不回傳；reduce_store_async 另有 pointer store overload，把結果以**弱寫入（weakly stored）**寫進 dst 指標。
  • inclusive_scan_update / exclusive_scan_update 除了原子更新外，會回傳「atomic 舊值與本 thread scan 結果結合後的值」（即該 thread 的偏移量，動態 buffer 配置即靠它取得各 thread 的 offset）。

val 限制：is_trivially_copyable == true，且 sizeof(T) <= 32（coalesced_group 與 ≤ 32 的 tile）或 sizeof(T) <= 8（更大 tile）；同一群組不同 thread 可傳不同 val。

sync / barrier_arrive / barrier_wait（sync.h）

arrive/wait 的 collective 約束

每個 phase，群組全員都必須各 arrive 一次、wait 一次。barrier_arrive 後到 barrier_wait 觀察到 phase 完成前，對該群組做任何 collective 或再次 arrive 皆 UB。barrier_wait 上的 thread 可能在其他 thread 還沒 wait 前就被釋放，但僅在全員都 arrive 之後。token 被 barrier_wait 消耗，不可重用。此機制讓 thread 在 arrive 後、wait 前做獨立工作以隱藏同步延遲。

Grid Synchronization

跨 grid 同步可讓應用改寫成 persistent thread blocks，免去 kernel 邊界的 state 失效與效能損失。

grid_group grid = this_grid();
grid.sync();

考試/測驗重點

題型 關鍵答案
__mbarrier_* 的 bar 必須位於哪種記憶體？ __shared__ memory
__mbarrier_arrive 與 __mbarrier_arrive_and_drop 差在哪？ 後者額外遞減下一 phase 的 expected count（drop 自身）
test_wait 與 try_wait 的差別？ test_wait 立即回傳查詢結果；try_wait 未完成時可能掛起，到 phase 完成或 max_sleep_nanosec 時限
*_parity 系列如何判定 phase？ true=奇數 phase、false=偶數 phase
重新利用 mbarrier 的 shared memory 前要做什麼？ 先 __mbarrier_inval(bar)
__pipeline_memcpy_async 的 size_and_align 合法值？ 只能是 4、8、16，且等於兩端指標的對齊
zfill 的作用？ 複製 size_and_align - zfill bytes，其餘補零
__pipeline_wait_prior(N) 等到哪些批次完成？ 提交索引 {0..L}，等到 L-N（含）
__pipeline_arrive_on 對 arrival count 的淨效果？ 0（先 +1，先前 memcpy_async 全完成後 −1）
取得三大 implicit group 的函式？ this_thread_block()／this_cluster()／this_grid()
cluster_group 需要什麼 compute capability？ CC 9.0+；非 cluster grid 視為 1×1×1
tiled_partition<Size> 對 Size 的限制？ 2 的次方且 ≤ 1024；父群大小須整除；native 1/2/4/8/16/32
shfl_up/shfl_down/shfl_xor/ballot 的 tile 尺寸限制？ 僅 ≤ 32
coalesced_threads() 的特性？ 機會性取得當下 active threads；meta_group_size=1、meta_group_rank=0
binary_partition 與 labeled_partition 關係？ binary 是 label 只能 0/1 的特例
cg::reduce 何時用硬體加速？ CC 8.0+ 且 add/min/max 或 AND/OR/XOR，且只有 4-byte 型別
inclusive vs exclusive scan？ inclusive 含呼叫者自身；exclusive 只含 thread_rank 較低者
cg::reduce / cg::scan 可用於哪些群組型別？ 僅 coalesced_group 與 thread_block_tile（不含 thread_block / grid_group）
cg::memcpy_async 兩種 overload 的計數單位差別？ shape 版以 bytes；dstLayout/srcLayout 版以元素（複製 min(dstLayout, srcLayout) 個）
*_update / *_async 變體用什麼做原子更新？ cuda::atomic/cuda::atomic_ref，relaxed memory ordering
barrier_arrive 回傳的 token 能重用嗎？ 不能，被 barrier_wait 消耗
grid.sync() 啟動 kernel 要用什麼 API？ cudaLaunchCooperativeKernel（非 <<<>>>）；CC 6.0+
查詢 cooperative launch 支援的 device attribute？ cudaDevAttrCooperativeLaunch