CUDA 中的 原子操作函数（atomic functions） 指能在并发线程中安全地对共享变量进行操作的函数。它们避免了竞态条件（race condition），保证操作的原子性（不可分割），即多个线程对同一个变量的读-改-写过程不会相互干扰。

所有原子函数都只能用于 __device__ 代码块中（即内核函数中）；

原子函数的返回值是操作执行前的原始值（即修改之前的值）。

常用原子操作函数分类

整型原子操作（整数类型 int, unsigned int, unsigned long long）

函数	说明
atomicAdd(address, val)	*address += val 原子加
atomicSub(address, val)	*address -= val 原子减
atomicExch(address, val)	*address = val 原子交换
atomicMin(address, val)	*address = min(*address, val)
atomicMax(address, val)	*address = max(*address, val)
atomicInc(address, limit)	(*address < limit) ? (*address)++ : 0
atomicDec(address, limit)	`(*address == 0
atomicCAS(address, compare, val)	比较并交换（compare-and-swap）
atomicAnd(address, val)	原子按位与
atomicOr(address, val)	原子按位或
atomicXor(address, val)	原子按位异或

浮点型原子操作（float, double）

在早期版本的 CUDA 中，浮点类型没有原子操作支持。但现在已经部分支持：

函数	说明	支持设备
atomicAdd(float* address, float val)	原子加	所有现代 GPU
atomicAdd(double* address, double val)	原子加	需要支持 Compute Capability >= 6.0（即 Pascal 架构及之后）

注意：CUDA 目前只支持浮点加法的原子操作，不支持 min/max/mul 等。

atomicCAS

atomicCAS(int* address, int compare, int val)

• 如果 *address == compare，则执行 *address = val；

• 否则不变。

• 返回原始值。

可以用它构建更复杂的原子逻辑（比如锁、更新最小值等）。

例子：使用 atomicCAS 实现一个简单自旋锁

__device__ void acquire_lock(int* lock) {
     // 尝试将 lock 从 0 设置为 1（即空闲变为占用）
    while (atomicCAS(lock, 0, 1) != 0); // 自旋直到成功获取
}

__device__ void release_lock(int* lock) {
    // 释放锁：把锁设为 0
    atomicExch(lock, 0); // 释放锁
}

支持 block 内共享锁 

你也可以把 lock 放在共享内存中，实现 block 内的线程同步：

__shared__ int block_lock;

if (threadIdx.x == 0) block_lock = 0;
__syncthreads();

acquire_lock(&block_lock);
// block 内线程临界区
release_lock(&block_lock);

地址空间支持

地址空间	是否支持原子操作
全局内存（global memory）	✅ 支持
共享内存（shared memory）	✅ 支持
本地内存（local memory）	❌ 不支持
常量内存（constant memory）	❌ 不支持
统一内存（unified memory）	✅ 支持，行为类似全局内存

使用示例代码啊

实现了统计字符流里的字符出现频率

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include <iostream>
#include<cstdio>
#define SIZE 100*1024
void error_handling(cudaError_t res) {
    if (res !=cudaSuccess) {
        std::cout << "error!" << std::endl;
    }
}
__global__ void histo_kernel(char* buffer, long size, unsigned int* histo) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd(&(histo[buffer[i]]), 1);
        i += stride;
    }
}
int main() {
    char* buffer=new char[SIZE];
    std::cin.getline(buffer, SIZE);
    unsigned int histo[256] = { 0 };//频率数组
    cudaEvent_t start, stop;
    error_handling(cudaEventCreate(&start));
    error_handling(cudaEventCreate(&stop));
    error_handling(cudaEventRecord(start, 0));
    //设备上的buffer和histo
    char* dev_buffer;
    unsigned int *dev_histo;
    error_handling(cudaMalloc((void**)&dev_buffer,SIZE));
    error_handling(cudaMalloc((void**)&dev_histo, 256*sizeof(int)));
    error_handling(cudaMemcpy(dev_buffer,buffer,SIZE,cudaMemcpyHostToDevice));
    error_handling(cudaMemset(dev_histo,0,256*sizeof(int)));//***

    cudaDeviceProp prop;
    error_handling(cudaGetDeviceProperties(&prop, 0));
    int blocks = prop.multiProcessorCount;
    int input_len = strlen(buffer);
    histo_kernel << <blocks * 2, 256 >> > (dev_buffer, input_len, dev_histo);

    error_handling(cudaEventRecord(stop, 0));
    error_handling(cudaEventSynchronize(stop));      // 等待 stop 事件代表的所有前序操作完成
    float elapsedTime;
    error_handling(cudaEventElapsedTime(&elapsedTime, start, stop));
    std::cout << "Kernel execution time: " << elapsedTime << " ms" << std::endl;

    error_handling(cudaMemcpy(histo,dev_histo,256*sizeof(int),cudaMemcpyDeviceToHost));
    
    for (int i = 0; i < 256; i++) {
        std::cout << static_cast<char>(i) << " occurred " << histo[i] << " times" << std::endl;
    }
    error_handling(cudaEventDestroy(start));
    error_handling(cudaEventDestroy(stop));
    error_handling(cudaFree(dev_buffer));
    error_handling(cudaFree(dev_histo));
    delete[]buffer;
}