同步原語與原子函式 (Synchronization Primitives and Atomic Functions)

重點總覽

CUDA 採用弱記憶體序模型 (weakly ordered memory model)：一個執行緒寫入的順序，不保證是另一個執行緒觀察到的順序。要安全地在執行緒間共享資料，必須靠三類工具：block/warp 同步原語（讓執行緒在某點集合並建立 happens-before）、memory fence（只排序、不保證可見性與原子性）、以及atomic 原子函式（不可分割的 read-modify-write）。三者解決的問題不同，常需組合使用。

Thread Block 同步：__syncthreads 家族

void __syncthreads();
int  __syncthreads_count(int predicate);
int  __syncthreads_and(int predicate);
int  __syncthreads_or(int predicate);

語意：

• 集合 (barrier)：等到 block 內所有未退出 (non-exited) 執行緒都到達同一個 __syncthreads*() 呼叫點，或已退出。
• 記憶體排序：呼叫 __syncthreads*() strongly happens-before（C++ 規範 [intro.races]）任何參與執行緒被解除等待或退出 — 因此同步前的共享/全域記憶體寫入，對解除等待後的執行緒都已排序完成。

帶 predicate 的變體（語意同 __syncthreads()，外加對所有未退出執行緒評估 predicate）：

// 假設 blockDim.x == 128
// 正確：條件對 block 內所有執行緒一致
if (blockIdx.x > 0) {        // CORRECT, uniform condition
    __syncthreads();
    output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
}

// 錯誤：條件依 threadIdx.x 而異 → 非 uniform
if (threadIdx.x > 0) {       // WRONG, non-uniform condition
    __syncthreads();         // Undefined Behavior / hang
}

// 錯誤：迴圈中 i == threadIdx.x 只對單一執行緒成立 → 非 uniform
for (int i = 0; i < blockDim.x; ++i) {
    if (i == threadIdx.x) {  // WRONG, non-uniform condition
        __syncthreads();     // Undefined Behavior
    }
}

典型用法：先各執行緒寫 __shared__ 陣列 → __syncthreads() → 由 thread 0 安全讀取整個陣列求和（保證所有寫入已排序在解除等待之前）。

Warp 同步：__syncwarp(mask)

void __syncwarp(unsigned mask = 0xFFFFFFFF);

• 協調同一 warp 內的執行緒通訊；用於避免同 warp 對共享/全域記憶體同址的 RAW / WAR / WAW 資料危害。
• 提供 mask 指名執行緒間的記憶體排序：__syncwarp(mask) strongly happens-before 任何 mask 指名的 warp 執行緒被解除等待或退出。
• 受 Warp __sync Intrinsic Constraints 約束（與其他 warp __sync intrinsics 相同的參與規則）。

if (threadIdx.x < warpSize) {
    __shared__ int shared_data[warpSize];
    shared_data[threadIdx.x] = input_data[threadIdx.x];
    __syncwarp();            // 等同 __syncwarp(0xFFFFFFFF)
    if (threadIdx.x == 0)
        output_data[0] = shared_data[1];
}

Memory Fence Functions（記憶體柵欄）

弱記憶體序下，若不加 fence 或同步就對同一位置讀寫，結果為未定義行為。Fence 強制記憶體存取呈現循序一致 (sequentially consistent) 的排序；差別只在於強制排序的 thread scope，而與存取的記憶體空間（shared / global / page-locked host / peer device）無關。

三層 scope 對照（intrinsic 與 libcu++ 等價寫法）：

語意（以 block-level 為例）：fence 前的所有寫入，被觀察者視為發生在 fence 後的所有寫入之前；fence 前的所有讀取也排序在 fence 後的讀取之前。

如何選 scope：

thread 1 寫 X,Y   ──fence──   thread 2 讀 Y,X
  同一 block            →  block-level fence 即可
  同 device 不同 block  →  device-level fence
  不同 device           →  system-level fence

無 fence 時，writeXY/readXY 對同址 X、Y 同時讀寫構成 data race（UB），A、B 可為任意組合（含 A=1,B=20）。加了 device-level fence 後，A=1,B=2 / A=10,B=20 / A=10,B=2 都可能發生，但 A=1,B=20 不可能——fence 保證對 X 的寫入可見於對 Y 的寫入之前。

跨 block 規約範例（單次 kernel 求 N 數總和）：每個 block 先算部分和並存入全域記憶體，再由最後完成的 block 把所有部分和加總。判斷誰最後：每個 block 以 atomicInc(&count, gridDim.x) 遞增計數器，最後一個 block 拿到的舊值等於 gridDim.x - 1。

Atomic Functions：五種提供方式

原子函式對共享資料做 read-modify-write，使其「要嘛全做、要嘛不做」，給所有參與執行緒一致的資料視圖。CUDA 提供五種：

Legacy 原子函式：運算、型別與 scope

對 global 或 shared memory 中的 32 / 64 / 128-bit word 做原子 RMW。

• 僅能用於 device 函式。
• 向量型別（__half2、__nv_bfloat162、float2、float4）：RMW 對每個元素分別保證原子，整個向量在單次存取中不保證原子。
• 記憶體序為 memory_order_relaxed，原子性只在特定 thread scope 成立：

各運算的計算式與型別支援：

atomicCAS 慣用法（萬用基底）

任何原子運算都可用 atomicCAS() (Compare-And-Swap) 迴圈實作，例如自製單精度 atomicAdd：

__device__ float customAtomicAdd(float* d_ptr, float value) {
    volatile unsigned* p = reinterpret_cast<unsigned*>(d_ptr);
    unsigned old_value = *p, assumed;
    do {
        assumed = old_value;
        float expected = cuda::std::bit_cast<float>(assumed) + value;
        unsigned expected_u = cuda::std::bit_cast<unsigned>(expected);
        // 用「整數比較」做 CAS，避免 NaN != NaN 造成永久 hang
        old_value = atomicCAS(p, assumed, expected_u);
    } while (assumed != old_value);
    return cuda::std::bit_cast<float>(old_value);
}

Built-in 原子函式 __nv_atomic_*

CUDA 12.8+ 提供，遵循 GNU atomic built-in 函式簽章並多一個 thread scope 參數；支援時 nvcc 定義巨集 __CUDACC_DEVICE_ATOMIC_BUILTINS__。

memory order 與 thread scope 列舉（須為整數字面量）：

order : __NV_ATOMIC_RELAXED, CONSUME, ACQUIRE, RELEASE, ACQ_REL, SEQ_CST
scope : __NV_THREAD_SCOPE_THREAD, BLOCK, CLUSTER, DEVICE, SYSTEM
        (預設參數通常為 __NV_THREAD_SCOPE_SYSTEM)

• __NV_ATOMIC_CONSUME 目前以較強的 __NV_ATOMIC_ACQUIRE 實作。
• __NV_THREAD_SCOPE_THREAD 目前以較寬的 __NV_THREAD_SCOPE_BLOCK 實作。
• __NV_THREAD_SCOPE_CLUSTER 需 sm_90 及更高架構。

限制：

• 僅能用於 device 函式；不可作用於 local memory。
• 不可取函式位址（不能當模板預設引數、不能放進建構子初始化列表等）。
• order 與 scope 引數必須是整數字面量，不能是變數。

運算族：fetch_* 回傳 old；無 fetch 的版本無回傳值。

考試/測驗重點

Related Notes

• 05-Technical-Appendices/06-Function-and-Variable-Annotations
• 05-Technical-Appendices/08-Warp-Functions-and-Macros
• 05-Technical-Appendices/12-Memory-Barrier-Pipeline-Cooperative-APIs
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 05-Technical-Appendices/14-CUDA-Cpp-Memory-Model
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps