CUDA C++ 記憶體模型 (CUDA C++ Memory Model)

重點總覽

標準 C++ 假設執行緒同步的成本是均勻且低的；CUDA C++ 不同：執行緒距離越遠，同步成本越高——block 內極低，跨多 GPU/CPU 的系統範圍則很高。為了表達這種非均勻成本，CUDA C++ 在 cuda:: namespace 中以 thread scope 擴充標準 C++ 記憶體模型與並行設施，預設仍保留標準 C++ 的語法與語意。

Thread Scopes

thread scope 指定「哪一類執行緒」可以透過某個同步原語（如 cuda::atomic、cuda::barrier）彼此同步。

namespace cuda {
enum thread_scope {
  thread_scope_system,
  thread_scope_device,
  thread_scope_block,
  thread_scope_thread
};
} // namespace cuda

Scope Relationships（涵蓋範圍）

每個程式執行緒透過一或多個 scope 關係彼此關聯，由大到小（外層涵蓋內層）：

system  ┌──────────────────────────────────────────┐ 多 GPU + CPU，成本最高
        │ device ┌───────────────────────────────┐ │ 單一 device（同 sync domain）
        │        │ block ┌──────────────────────┐ │ │ 單一 thread block，成本最低
        │        │       │ thread (僅自身)      │ │ │
        │        │       └──────────────────────┘ │ │
        │        └───────────────────────────────┘ │
        └──────────────────────────────────────────┘

Synchronization primitives

std:: 與 cuda::std:: 中的型別，當以 cuda::thread_scope_system 實例化對應的 cuda:: 型別時，行為相同。

• 換言之，標準 C++ 的同步型別（如 std::atomic）等同於最大範圍（system scope）的 cuda:: 版本。
• 要取得 block/device 等較窄、成本較低的 scope，需改用 cuda::atomic、cuda::atomic_ref、cuda::barrier 等並以 scope 作為模板參數。

Atomicity

一個 atomic 操作只在它所指定的 scope 下保證為 atomic，條件為以下之一：

• 它指定的 scope 不是 cuda::thread_scope_system；或
• scope 是 cuda::thread_scope_system 且滿足下列任一：

GPU memory 且僅 GPU 執行緒存取時，需再滿足任一：

• 對每個存取的 srcDev 與物件所在 GPU dstDev，cudaDeviceGetP2PAttribute(&val, cudaDevP2PAttrNativeAtomicSupported, srcDev, dstDev) 為 1；或
• 只有來自單一 GPU 的 GPU 執行緒並行存取它。

Data Races

CUDA C++ 對 ISO/IEC 14882（C++ 標準）intro.races 第 21 段做如下修改（加入 scope 概念）：

• 關鍵：是否 atomic 要看「該操作指定的 scope 是否涵蓋對方執行緒」，而非單看操作本身是否為 atomic 指令。
• barrier、latch、counting_semaphore 的成員函式（解構子除外）並行呼叫不會引入 data race，如同它們是 atomic 操作。
• stop_token 的 request_stop/stop_requested/stop_possible 這三個特定函式的呼叫也不會引入 data race（來源僅針對這三個函式，且未含「解構子除外」的限定）。
• Fence 同步（atomics.fences）也加上 scope 約束，分三種 synchronizes-with 配對：
  • release fence A 與 acquire fence B：存在 atomic operations X、Y，**四個操作（A、B、X、Y）**各自都必須指定涵蓋其他操作之執行緒的 scope。
  • release fence A 與執行 acquire 的 atomic operation B：只牽涉 (A、B、X) 三個操作，各自須指定涵蓋對方執行緒的 scope。
  • release operation A 與 acquire fence B：同樣只牽涉 (A、B、X) 三個操作，各自須指定涵蓋對方執行緒的 scope。

Example: Message Passing

Block 0 的執行緒把訊息寫進 x，透過旗標 f 傳給 Block 1 的執行緒。正確版：兩端都用 thread_scope_device（涵蓋跨 block 的兩個執行緒）。

// 初始
int x = 0, f = 0;

// Thread 0, Block 0
x = 42;
cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
flag.store(1, memory_order_release);

// Thread 0, Block 1
cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
while (flag.load(memory_order_acquire) != 1);
assert(x == 42); // 必成立：release/acquire 在 device scope 下成對同步

錯誤版（data race / UB）：store 端誤用 thread_scope_block。

// Thread 0, Block 0
x = 42;
cuda::atomic_ref<int, cuda::thread_scope_block> flag(f);
flag.store(1, memory_order_release); // UB: data race

// Thread 0, Block 1
cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
while (flag.load(memory_order_acquire) != 1); // UB: data race
assert(x == 42);

考試/測驗重點

Related Notes

• 05-Technical-Appendices/15-CUDA-Cpp-Execution-Model
• 05-Technical-Appendices/07-Synchronization-and-Atomic-Functions
• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps