Thread Scopes 與 Scoped Atomics (Thread Scopes and Scoped Atomics)

重點總覽

3.2.3 Thread Scopes（執行緒範圍）

CUDA 執行緒構成 Thread Hierarchy；在此階層中，記憶體操作的「可見性」與「同步範圍」並非一致。為描述這種非一致性，CUDA 程式設計模型引入 thread scope 的概念。

• thread scope 的定義：決定哪些執行緒能觀察到某執行緒的 loads 與 stores，並指定哪些執行緒能透過 atomic operations、barriers 等 synchronization primitives 互相同步。
• 每個 scope 都對應記憶體階層中的一個 point of coherency（一致性點）。
• thread scope 暴露於 CUDA PTX，同時以擴充形式提供於 libcu++ 函式庫（cuda::thread_scope_*）。

scope 涵蓋範圍與一致性點由窄到寬：

 scope               PTX        coherency point        可見對象
 thread          (private)        –                    僅自己
   │
 block            .cta            L1                    同 block 的執行緒
   │
 cluster          .cluster        L2                    同 cluster 的執行緒
   │
 device           .gpu            L2                    同 GPU 的執行緒
   │
 system           .sys            L2 + connected caches CPU + 其他 GPU
   ▼  (越往下涵蓋越廣、需穿越的快取階層越多 → 越慢)

3.2.4 Advanced Synchronization Primitives 概觀

本節介紹三大類同步原語：

本筆記聚焦 Scoped Atomics；Asynchronous Barriers 與 Pipelines 見 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines。

3.2.4.1 Scoped Atomics

Scoped atomics 支援 C++ 標準 atomic 記憶體語意，可透過 libcu++ 函式庫或 編譯器 built-in functions 使用。它在 CUDA thread hierarchy 的適當層級提供高效同步，兼顧正確性與效能。

兩大組成（3.2.4.1.1）

Scoped atomics 結合兩個關鍵概念：

• Thread Scope：定義哪些執行緒能觀察到此 atomic 操作的效果。
• Memory Ordering：定義此操作相對其他記憶體操作的順序約束。

範例一：block-scoped 原子計數器

#include <cuda/atomic>
__global__ void block_scoped_counter() {
    // 只在本 block 內可見的 shared atomic counter
    __shared__ cuda::atomic<int, cuda::thread_scope_block> counter;
    if (threadIdx.x == 0) {                       // 只由 thread 0 初始化
        counter.store(0, cuda::memory_order_relaxed);
    }
    __syncthreads();                              // 確保所有執行緒看到初始化結果
    int old_value = counter.fetch_add(1, cuda::memory_order_relaxed);
    // 使用 old_value...
}

等價的 built-in 寫法以 __nv_atomic_* 搭配 __NV_ATOMIC_RELAXED、__NV_THREAD_SCOPE_BLOCK：

__shared__ int counter;
if (threadIdx.x == 0) {
    __nv_atomic_store_n(&counter, 0, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_BLOCK);
}
__syncthreads();
int old_value = __nv_atomic_fetch_add(&counter, 1,
                                      __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_BLOCK);

此例展示 scoped atomics 的基本要點：

• Shared Variable：以 __shared__ 記憶體在 block 內共享單一 counter。
• Atomic Type Declaration：cuda::atomic<int, cuda::thread_scope_block> 建立具 block 層級可見性的 atomic 整數。
• Single Initialization：只由 thread 0 初始化，避免設定期間的 race condition。
• Block Synchronization：__syncthreads() 確保所有執行緒在繼續前都看到已初始化的 counter。
• Atomic Increment：每個執行緒原子地遞增 counter 並取回先前值。

範例二：producer-consumer 用 acquire-release 確保順序

__global__ void producer_consumer() {
    __shared__ int data;
    __shared__ cuda::atomic<bool, cuda::thread_scope_block> ready;
    if (threadIdx.x == 0) {
        // Producer：先寫 data，再 signal ready
        data = 42;
        ready.store(true, cuda::memory_order_release);  // release 確保 data 寫入可見
    } else {
        // Consumer：等 ready 訊號後再讀 data
        while memory_order_acquire) { // acquire 確保讀到該寫入
            // spin wait
        }
        int value = data;
        // 處理 value...
    }
}

producer-consumer 的 release/acquire 時序：

 Producer (thread 0)                 Consumer (thread != 0)
 ───────────────────                 ──────────────────────
 data = 42;                          while(!ready.load(acquire)) { spin }
 ready.store(true, release) ─────────►          │
        │   release: 之前的寫入                  │  acquire: 之後的讀取
        │   不能重排到 store 之後                 │  不能重排到 load 之前
        ▼                                        ▼
   (data 寫入對 consumer 可見) ───── 配對 ───► int value = data;  // 保證讀到 42

Memory Ordering 語意對照（C++ 標準）

3.2.4.1.2 Performance Considerations（效能考量）

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines
• 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config
• 03-Advanced-CUDA/09-Multi-GPU-Programming
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps