一、先明确：CUDA原子操作本质是线程级的 #

所有CUDA原子操作都是单个线程发起的操作——核心是保证"读-改-写"这一整套动作不会被其他线程打断，同一时间只有发起操作的那个线程能完成这个完整流程。不管是操作全局内存还是共享内存，原子操作的执行主体都是线程，所以本质上是线程级的。

二、搞懂"块级序列化"到底是什么 #

你说的"块级序列化"，其实是一种跨线程块的执行阻塞现象：当多个不同线程块里的线程，同时对同一个全局内存地址执行原子操作时，GPU的全局原子处理单元没办法同时处理这些请求，只能让这些线程块里的线程排队执行——也就是一个块的线程先完成原子操作，下一个块的线程才能跟上，原本的多块并行就变成了串行，这就是所谓的"块级序列化"。

举个例子：如果10个线程块都往同一个全局变量total_sum里做atomicAdd，那这10个块的原子操作就会排队，一个接一个来，没办法并行推进，这就是块级序列化。

而如果是同一个线程块内的线程对共享内存做原子操作，因为共享内存是每个块私有的，而且SM（流多处理器）有本地的原子处理单元，同一SM内的线程操作共享内存原子时，只会有线程级的排队，不会影响其他块的并行，也就不会触发块级序列化。

三、什么时候会出现块级序列化的原子操作？ #

主要有两种典型场景：

• 多个线程块的线程同时操作同一个全局内存地址（比如多块累加同一个全局变量）；
• 全局原子单元成为性能瓶颈，来自不同块的原子请求积压，只能串行处理，此时就表现为块级的序列化。

四、你的代码会不会触发块级序列化？ #

先看你的代码（顺便提个小问题，共享内存的初始化写法不对，后面会纠正）：

__global__ void sum (int *input){ 
    if ( threadIdx.x == 0) 
        __shared__ int result = 0; 
    __syncthreads(); 
    atomicAdd(result,input[threadIdx.x+blockDim.x*blockId.x]); 
}

首先纠正个小错误：__shared__变量不能在if分支里声明初始化，正确的写法应该是先声明共享变量，再由线程0初始化后同步：

__shared__ int result;
if (threadIdx.x == 0) result = 0;
__syncthreads();

回到块级序列化的问题：这里的atomicAdd是对**当前线程块私有的共享内存变量result**执行的。每个线程块的result都是独立的，不同块的线程根本不会碰彼此的result，所以完全没有跨块的内存竞争。

结论是：这段代码不会发生块级序列化。各个线程块的原子操作各自在自己的共享内存里完成，彼此并行不干扰，只有同一个块内的线程会有线程级的原子排队，但这和块级序列化完全是两回事。

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