CUDA 原子操作线程或块级别？

Question

什么时候发生在块级别序列化的原子操作？如果我有以下代码：

__global__ void sum (int *input){

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

此序列化是否发生在块级别？我一般不明白 "block level serialization" 是什么意思，因为据我所知，操作总是由线程执行。

Answer 1

CUDA中有两种原子操作：

那些对全局内存进行操作的
那些在共享内存上操作的

全局内存对grid/kernel中的所有线程都是"visible"（即全局内存只有一个逻辑视图，网格中的所有线程共享相同的视图），因此是全局的原子创建（根据需要）序列化，即 device-wide，通常 (*) 在 L2 缓存中解析，这是一个 device-wide 资源。

共享内存在逻辑上是一个 per-threadblock 资源（不同的线程块对共享内存有不同的逻辑视图，即它们自己的 "private copy" 共享内存）并且在物理上是一个资源 per-SM.由于共享内存在逻辑上是一种 per-threadblock 资源，因此只有特定线程块中的线程具有相同的 "view" 或共享内存的逻辑副本。因此，在执行 shared-memory 原子时，可能发生的 "serialization" 是线程块级别的序列化（与设备范围相比）。只有同一线程块中的线程才能竞争访问特定的共享内存位置，因为共享内存在逻辑上是每个线程块的独立实体。

在单线程内，所有的指令都是逻辑序列化的，没有只在单线程内运行的原子操作的概念。原子操作要么（逻辑上）跨属于单个线程块的线程，要么跨 grid/kernel.

中的所有线程。

当然，所有原子操作都涉及线程之间为访问特定位置而进行的竞争。重要的区别在于它们是在竞争共享内存中的某个位置，还是全局内存中的某个位置。

在语法上，共享内存原子和全局内存原子之间没有区别（在 CUDA C++ 中，即；CUDA PTX 或 SASS 是不同的情况）。确定原子操作是在共享内存还是全局内存上发生是传递给提供更新位置的原子的指针的类型（即数值，或者逻辑关联）的函数。如果那个指针 "points" 指向共享内存，它就是一个共享内存原子。如果那个指针"points"指向全局内存，它就是一个全局内存原子。

原子是documented

(*) 最近，nvcc 编译器驱动程序获得了识别某些原子模式 warp-wide 并执行 atomic aggregation 的能力。当编译器选择这个实现时，可以说当编译器选择那个惯用语时，warp 内的原子竞争没有在 L2 缓存中得到解决。

CUDA 原子操作线程或块级别？

CUDA Atomic Operations thread or block level?

cuda

atomic