3. 同步、原子与内存模型

CPU 程序员有三件套：mutex、条件变量、std::atomic。GPU 上只剩最后一件，而且语义更弱。这一节讲清楚 __syncthreads() 到底保证了什么、原子什么时候会变成瓶颈，以及为什么“跨 block 自旋等待”是个陷阱。

3.1 block 内的 barrier

__syncthreads() 做两件事：所有线程到达后才继续，以及此前的 shared/全局写对 block 内其他线程可见。它相当于“barrier + block 范围的 fence”。

exchange.cu

__shared__ float s[256];
int t = threadIdx.x;
s[t] = x[i];
__syncthreads();                 // 少这一行，读到的可能是旧值
float left = s[(t + 255) & 255]; // 取左邻居，边界回绕

两个硬性限制：

• 只能同步 block 内。不同 block 之间没有 barrier；__syncthreads() 不会让 block 1 等 block 0。
• 必须在所有线程都到达的位置。如果写在 if (threadIdx.x < 16) { … __syncthreads(); } 里，有的 lane 到了、有的没到，行为是未定义（常见表现是挂死或结果随机）。条件必须块内一致。

还有一个容易忽略的点：__syncthreads() 之后再次写同一块 shared 之前，往往需要第二个 __syncthreads()——否则快线程可能把慢线程还没读的数据覆盖掉。归约、转置、双缓冲都踩过这个坑。

3.2 原子操作：能用，但要少用

atomicAdd、atomicSub、atomicMin/Max、atomicExch、atomicCAS、按位与/或/异或，都返回旧值。它们是设备范围的，落在 L2 上完成。

count.cu

__global__ void count_positive(const float* x, int n, int* counter) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n && x[i] > 0.0f)
        atomicAdd(counter, 1);
}

这段代码是对的，但很慢：如果数据里正数很多，所有 block 都在抢同一个地址，L2 上的一次原子要串行几十到几百个周期。正确做法是先在 block 内归约，再让每个 block 做一次原子（4.1）；或者用 shared 私有化（4.4）。

原子还有第二个代价：不可复现。到达顺序由调度决定，而浮点加法不满足结合律——同一份输入跑两次，结果可能差几个 ULP（机制在 3.3 末尾）。所以不要把「汇总」交给原子：要可复现，就换成固定树形，写法在 4.1 末尾。

3.3 弱内存模型与 fence

GPU 的内存模型是弱（relaxed）的：没有 fence 时，编译器和硬件都可以重排你的读写。下面这段“发布数据”的代码，如果去掉 __threadfence()，其他 SM 可能先看到 g_ready == 1，再看到旧的 *data：

publish.cu

__device__ int g_ready = 0;

__global__ void publish(float* data) {
    if (blockIdx.x == 0 && threadIdx.x == 0) {
        *data = 1.0f;
        __threadfence();            // 先把 data 的写推到设备可见
        atomicExch(&g_ready, 1);    // 再发布标志
    }
}

消费端的顺序相反：先读标志，再读数据，并且需要 acquire 语义。干净的写法是用 libcu++ 的显式内存序：

#include <cuda/atomic>

cuda::atomic_ref<int, cuda::thread_scope_device> flag(g_ready);
while (flag.load(cuda::memory_order_acquire) == 0) { }
float v = *data;   // 现在能保证看到 publish 写入的值

三个 fence 的范围：__threadfence_block()（block 内）、__threadfence()（设备内）、__threadfence_system()（含主机）。它们只做排序，不做互斥。

这也是为什么“原子操作的结果不可复现”：atomicAdd(float) 的到达顺序由调度决定，浮点加法不满足结合律，每次运行的舍入都不同。CPU 多线程累加也有同样的问题，只是 GPU 上更容易暴露。

3.4 warp 级原语

warp 内的 32 个 lane 可以直接交换寄存器，不需要 shared memory，也不需要 __syncthreads()：

__inline__ __device__ float warpSum(float v) {
    #pragma unroll
    for (int off = 16; off > 0; off >>= 1)
        v += __shfl_down_sync(0xffffffffu, v, off);  // lane i 取 i+off 的值
    return v;                                        // lane 0 得到总和
}

常用的一组：

mask 必须是当前活跃 lane 的集合。标准写法是全 1，但那要求 32 个 lane 都在执行这条指令。如果 blockDim.x 不是 32 的倍数，最后一个 warp 只有部分 lane 活跃，全 1 的 mask 就是错的——这是 1.7 里提到的那个坑。

unsigned m = __ballot_sync(0xffffffffu, v > 0.0f);
int leader = __ffs(m) - 1;          // 最低位满足条件的 lane
int npos  = __popc(m);              // 满足条件的 lane 数

3.5 和 CPU 的对照

3.6 容易写错的地方