非同步 Barriers 與 Pipelines (Asynchronous Barriers and Pipelines)

重點總覽

Asynchronous Barriers（非同步 barrier）

Asynchronous barrier 與典型的單階段 barrier（__syncthreads()）不同：thread 發出「我到了」的通知（arrival）被從「等待其他 thread 到達」的動作（wait）中分離出來。這個分離提升執行效率，因為 thread 可在等待期間執行與 barrier 無關的額外運算，更有效利用等待時間。

關鍵用途：

• 用 CUDA threads 實作 producer-consumer 模式。
• 讓記憶體階層內的非同步資料複製在完成時對 barrier 發出訊號（「arrive on」barrier）。

硬體與可用性

• 在 compute capability 7.0 或更高的裝置上可用。
• CC 8.0+ 對 shared-memory 的 async barrier 提供硬體加速，且大幅提升同步粒度：可硬體加速地同步 block 內任意 subset 的 CUDA threads。
• 先前架構只能加速 whole-warp（__syncwarp()）或 whole-block（__syncthreads()）層級的同步。

API 層級

• CUDA 透過 cuda::std::barrier（libcu++ 中符合 ISO C++ 的 barrier）提供 async barrier；常用別名為 cuda::barrier。
• 除實作 std::barrier 外，函式庫提供 CUDA 專屬擴充，可選擇 barrier 的 thread scope 以改善效能，並暴露較低階的 cuda::ptx API。
• cuda::barrier 可透過 friend function cuda::device::barrier_native_handle() 取得 native handle，傳給 cuda::ptx 函式而互通。
• CUDA 另提供一組 primitives API，針對 thread-block scope 的 shared-memory async barrier。

不同 thread scope 的 async barrier 對照

同步的時間性切分（Temporal Splitting）

沒有 arrive-wait barrier 時，thread block 內同步靠 __syncthreads() 或 Cooperative Groups 的 block.sync()：

#include <cooperative_groups.h>
__global__ void simple_sync(int iteration_count) {
    auto block = cooperative_groups::this_thread_block();
    for (int i = 0; i < iteration_count; ++i) {
        /* code before arrive */
        // Wait for all threads to arrive here.
        block.sync();
        /* code after wait */
    }
}

threads 被阻擋在同步點直到所有 thread 都到達；且同步點之前的記憶體更新保證在同步點之後對 block 內所有 thread 可見。此模式有三階段：sync 前的更新 → 同步點 → sync 後（可見前面的更新）。

改用 async barrier 後，同步被切成 arrive 點與 wait 點：

#include <cuda/barrier>
#include <cooperative_groups.h>
__global__ void split_arrive_wait(int iteration_count, float *data) {
    using barrier_t = cuda::barrier<cuda::thread_scope_block>;
    __shared__ barrier_t bar;
    auto block = cooperative_groups::this_thread_block();
    if (block.thread_rank() == 0) {
        init(&bar, block.size());   // expected arrival count = block 大小
    }
    block.sync();
    for (int i = 0; i < iteration_count; ++i) {
        /* code before arrive */
        barrier_t::arrival_token token = bar.arrive();  // arrive 不阻擋 thread
        compute(data, i);                                // 與等待重疊的獨立工作
        bar.waitmove(token);                      // 在此阻擋直到全員 arrive
        /* code after wait */
    }
}

重點：barrier 用 init(&bar, count) 設定預期到達數 (expected arrival count)（通常為 block.size()），且只由一個 thread 初始化再 block.sync()。bar.arrive() 回傳 arrival_token；中間的 compute() 重疊掉等待時間；bar.waitmove(token) 才真正阻擋。

token 與 barrier phase

• bar.arrive() 回傳的 arrival_token 擷取了 barrier 目前的 phase。
• bar.waitmove(token)) 會阻擋，直到參與 threads 完成 bar.arrive( 達到預期次數（即該 phase 完成）。
• phase 完成後 barrier 自動進入下一個 phase 並重設計數，因此可在迴圈中重複使用同一個 barrier。

記憶體可見性與五階段模型

arrive() 之前發生的記憶體更新，保證在參與 threads 呼叫 wait() 之後對它們可見。arrive() 不會阻擋 thread——它可以接著做不依賴其他 thread arrive 前更新的工作。

arrive-wait 模式的五階段：

Thread A                         Thread B
────────                         ────────
code before arrive               code before arrive
   │  (更新將被 wait 後讀取)         │
 arrive() ─┐  token              arrive() ─┐  token
   │       │ (非阻擋, 隱含 fence)    │       │
compute()  │  ← 重疊區 →          compute()  │
   │       │                       │       │
 wait(tok)◄┘                     wait(tok)◄┘  (阻擋直到全員 arrive)
   │  (可見對方 arrive 前更新)        │
code after wait                  code after wait

較低階 API：cuda::ptx 與 primitives

三套 API 是同一模式的不同層級寫法，差異在於 wait 的呼叫方式：

// cuda::ptx：手動輪詢 try_wait
cuda::ptx::mbarrier_init(&bar, block.size());          // bar 為 __shared__ uint64_t
uint64_t token = cuda::ptx::mbarrier_arrive(&bar);
while mbarrier_try_wait(&bar, token) {}  // 自旋等待

// C primitives：__mbarrier_*，try_wait 帶逾時(ns)上限
__mbarrier_init(&bar, block.size());                   // bar 為 __shared__ __mbarrier_t
__mbarrier_token_t token = __mbarrier_arrive(&bar);
while (!__mbarrier_try_wait(&bar, token, 1000)) {}

重點：cuda::ptx 與 primitives 用 try_wait 迴圈取代 bar.wait()；primitives 的 token 型別為 __mbarrier_token_t，且 __mbarrier_try_wait 第三參數是逾時上限。三者語意相同。

Producer-consumer 與 warp specialization

由於 CC 8.0+ 可硬體加速地同步 block 內任意 subset 的 threads，async barrier 很適合 warp specialization：把不同 warp（thread subset）指派成 producer 或 consumer，各自只在自己的 barrier 上 arrive/wait，彼此用 barrier 的 arrive 訊號協調，而非全 block 一起 __syncthreads()。

warp 0 (producer)            warp 1 (consumer)
─────────────────            ─────────────────
load → shared buffer
 arrive(full) ──────────────► wait(full)
       │ (繼續預取下一批)         │  consume shared buffer
 wait(empty) ◄────────────── arrive(empty)
 (buffer 釋放後再 load)

Pipelines（pipeline 同步物件）

CUDA 提供 pipeline 同步物件作為協調機制，把非同步 memory copy 排序成多個階段 (stages)，方便實作 double-buffering / multi-buffering 的 producer-consumer 模式。

• pipeline 是一個有 head 與 tail 的雙端佇列 (deque)，以 **FIFO（先進先出）**順序處理工作。
• producer threads 把工作 commit 到 head；consumer threads 從 tail 取出工作。
• 透過 libcu++ 的 cuda::pipeline API 與一組 primitives API 暴露。

     producer 端 (head)                          consumer 端 (tail)
  producer_acquire ─► 發 memcpy_async ─► producer_commit
        │                                                │
        ▼            FIFO ring of stages                 ▼
   ┌──────┬──────┬──────┬──────┐
   │ s0   │ s1   │ s2   │ s3   │   ◄ 多 buffer 同時在飛
   └──────┴──────┴──────┴──────┘
        ▲                                                │
        └──── consumer_release ◄─ 用 shared 資料 ◄─ consumer_wait

cuda::pipeline API

Primitives API

兩套 API 的比較

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps