CUDA内存模型：为何无需获取栅栏防止加载-加载重排序？

我正在阅读《Programming Massively Parallel Processors》一书，注意到以下实现“多米诺式”扫描的代码片段：

if (threadIdx.x == 0) {
    while(AtomicAdd(&flags[bid], 0) == 0) {}
    // ???? why do I not need thread fence here (for acquire semantic) to prevent load-load reordering between the loads for flags[bid] and scan_value[bid]? 
    previous_sum = scan_value[bid];
    scan_value[bid+1] = previous_sum + local_sum;
    __threadfence(); // why the "release" fence here is sufficient?
    atomicAdd(&flags[bid + 1], 1);
}

书中指出，__threadfence()需用于两次写入之间，确保scan_value[bid+1]在flags[bid+1]递增前写入全局内存，这类似C++内存模型的释放语义（__threadfence具备seq_cst语义）。但我无法理解：为何两次读取间无需类似的__threadfence（或获取栅栏）来防止加载-加载重排序（即先读取scan_value[bid]再读取flags[bid]）？

解答 #

1. 原子操作自带隐式获取语义 #

你用的AtomicAdd(&flags[bid], 0)本质是原子加载操作（加0不修改值，仅完成原子读取）。在CUDA内存模型中，所有原子操作本身就具备获取语义：原子加载执行完成后，后续的所有内存加载、存储操作都不能被重排到这个原子加载之前。

也就是说，previous_sum = scan_value[bid]这行读取操作，绝对不会被重排到AtomicAdd(...,0)之前执行——硬件和编译器都会严格保证这个顺序，不需要额外添加栅栏。

2. 加载重排的约束限制 #

CUDA架构对加载操作的重排有严格约束：只有当两个加载操作无内存地址依赖、且不涉及原子操作或栅栏时，才可能被重排。但这里的原子加载已经强制了执行顺序：线程必须先确认flags[bid]不为0，才能去读取scan_value[bid]，这种逻辑依赖加上原子操作的语义，直接阻断了重排的可能。

3. 释放栅栏的核心作用 #

后面的__threadfence()是为了保证scan_value[bid+1]的写入操作全局可见后，再去更新flags[bid+1]。atomicAdd(&flags[bid+1],1)属于原子存储操作，__threadfence()的seq_cst语义会确保之前的所有写入操作完成并同步到全局内存后，才执行这个原子递增，让下一个线程能读取到正确的scan_value值。

内容的提问来源于stack exchange，提问作者Lifu Huang