Atomic Functions

原子函数对驻留在全局或共享内存中的一个 32 位或 64 位字执行读取-修改-写入原子操作。 在 float2 或 float4 的情况下，对驻留在全局内存中的向量的每个元素执行读取-修改-写入操作。 例如，atomicAdd() 在全局或共享内存中的某个地址读取一个字，向其添加一个数字，然后将结果写回同一地址。 原子函数只能在设备函数中使用。

本节中描述的原子函数具有排序 cuda::memory_order_relaxed 并且仅在特定范围内是原子的：

  • 具有 _system 后缀的原子 API（示例：__atomicAdd_system）在 cuda::thread_scope_system 范围内是原子的。

  • 没有后缀的原子 API（例如：__atomicAdd）在 cuda::thread_scope_device 范围内是原子的。

  • 带有 _block 后缀的原子 API（例如：__atomicAdd_block）在 cuda::thread_scope_block 范围内是原子的。

在以下示例中，CPU 和 GPU 都自动更新地址 addr 处的整数值：

__global__ void mykernel(int *addr) {
  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
}

void foo() {
  int *addr;
  cudaMallocManaged(&addr, 4);
  *addr = 0;

   mykernel<<<...>>>(addr);
   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
}

注意，任何原子操作都可以基于 atomicCAS() (Compare And Swap) 来实现。 例如，双精度浮点数的 atomicAdd() 在计算能力低于 6.0 的设备上不可用，但可以按如下方式实现：
查看计算能力：https://developer.nvidia.com/zh-cn/cuda-gpus#compute

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif

以下设备范围的原子 API 有系统范围和块范围的变体，但以下情况除外：

  • 计算能力小于 6.0 的设备只支持设备范围的原子操作，

  • 计算能力低于 7.2 的 Tegra 设备不支持系统范围的原子操作。

算术函数

atomicAdd()

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address, unsigned long long int val);

读取位于全局或共享内存中address的 16 位、32 位或 64 位 old，计算 (old + val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);

atomicAdd() 的 32 位浮点版本仅受计算能力 2.x 及更高版本的设备支持。
atomicAdd() 的 64 位浮点版本仅受计算能力 6.x 及更高版本的设备支持。

__half2 atomicAdd(__half2 *address, __half2 val);
__half atomicAdd(__half *address, __half val);
__nv_bfloat162 atomicAdd(__nv_bfloat162 *address, __nv_bfloat162 val);
__nv_bfloat16 atomicAdd(__nv_bfloat16 *address, __nv_bfloat16 val);

atomicAdd() 的 32 位 half2 浮点版本仅受计算能力 6.x 及更高版本的设备支持。 对于两个 half 或 nv_bfloat16 元素中的每一个，分别保证 half2 或 nv_bfloat162 添加操作的原子性； 作为单个 32 位访问，不能保证整个 half2 或 __nv_bfloat162 是原子的。

atomicAdd() 的 16 位 __half 浮点版本仅受计算能力 7.x 及更高版本的设备支持。

atomicAdd() 的 16 位 __nv_bfloat16 浮点版本仅受计算能力 8.x 及更高版本的设备支持。

float2 atomicAdd(float2* address, float2 val);
float4 atomicAdd(float4* address, float4 val);

atomicAdd() 的 float2 和 float4 浮点向量版本仅受计算能力 9.x 及更高版本的设备支持。 float2 或 float4 add 操作的原子性分别为两个或四个 float 元素中的每一个保证； 作为单个 64 位或 128 位访问，不能保证整个 float2 或 float4 是原子的。

atomicAdd() 的 float2 和 float4 浮点向量版本仅受计算能力 9.x 及更高版本的设备支持。
atomicAdd() 的 float2 和 float4 浮点向量版本仅支持全局内存地址。

atomicSub()

int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address, unsigned int val);

读取位于全局或共享内存中地址地址的 32 位字 old，计算 (old - val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicExch()

int atomicExch(int* address, int val);
unsigned int atomicExch(unsigned int* address,　unsigned int val);
unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val);
float atomicExch(float* address, float val);

读取位于全局或共享内存地址处的 32 位或 64 位字 old，并将 val 存储回同一地址的内存。 这两个操作在一个原子事务中执行。 该函数返回旧的。

atomicMin()

int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address, unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address, unsigned long long int val);
long long int atomicMin(long long int* address, long long int val);

读取位于全局或共享内存地址的 32 位或 64 位字 old，计算 old 和 val 的最小值，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

64 位版本的 atomicMin() 仅受计算能力 5.0 及更高版本的设备支持。

atomicMax()

int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address, unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address, unsigned long long int val);
long long int atomicMax(long long int* address, long long int val);

读取位于全局或共享内存地址处的 32 位或 64 位字 old，计算 old 和 val 的最大值，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

64 位版本的 atomicMax() 仅受计算能力 5.0 及更高版本的设备支持。

atomicInc()

unsigned int atomicInc(unsigned int* address, unsigned int val);

读取位于全局或共享内存中地址地址的 32 位字 old，计算 (((old == 0) || (old > val)) ? val : (old-1) )，并将结果存回 到同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicCAS()

int atomicCAS(int* address, int compare, int val);
unsigned int atomicCAS(unsigned int* address,
                       unsigned int compare,
                       unsigned int val);
unsigned long long int atomicCAS(unsigned long long int* address,
                                 unsigned long long int compare,
                                 unsigned long long int val);
unsigned short int atomicCAS(unsigned short int *address,
                             unsigned short int compare,
                             unsigned short int val);

读取位于全局或共享内存中地址地址的 16 位、32 位或 64 位字 old，计算 (old == compare ? val : old)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的（比较和交换）。

Bitwise Functions

atomicAnd()

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址address 的32 位或64 位字old，计算(old & val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicAnd() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

atomicOr()

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
                      unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
                                unsigned long long int val);

读取位于全局或共享内存中地址地址的 32 位或 64 位字 old，计算 (old | val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicOr() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

atomicXor()

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
                                 unsigned long long int val);

读取位于全局或共享内存中地址地址的 32 位或 64 位字 old，计算 (old ^ val)，并将结果存储回同一地址的内存。 这三个操作在一个原子事务中执行。 该函数返回旧的。

atomicXor() 的 64 位版本仅受计算能力 5.0 及更高版本的设备支持。

参考：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#bitwise-functions