// 整数类型atomicAdd(int* address, int val);        // 加法atomicSub(int* address, int val);        // 减法atomicExch(int* address, int val);       // 交换atomicMin(int* address, int val);        // 最小值atomicMax(int* address, int val);        // 最大值atomicInc(int* address, int val);        // 增加到val后归0atomicDec(int* address, int val);        // 减少到0后重置为valatomicCAS(int* address, int compare, int val);  // 比较并交换// 64位整数（需要计算能力3.5+）atomicAdd(long long* address, long long val);// 浮点数（需要计算能力6.0+）atomicAdd(float* address, float val);atomicAdd(double* address, double val);加法示例__global__ void atomicAddKernel(int* counter, int* data, int N) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;        if (idx < N) {        // 每个线程将data[idx]的值原子性地加到counter        atomicAdd(counter, data[idx]);    }}计算直方图__global__ void histogramKernel(int* input, int* histogram, int N, int bins) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        int bin = input[idx] % bins;  // 计算bin索引        atomicAdd(&histogram[bin], 1);  // 原子递增对应bin    }}寻找最大值__global__ void findMaxKernel(int* input, int* maxValue, int N) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        int old = *maxValue;        int newVal = input[idx];        // 原子比较并更新最大值        while (newVal > old) {            if (atomicCAS(maxValue, old, newVal) == old) {                break;            }            old = *maxValue;  // 如果失败，重试        }    }}

CAS 流程

2.位运算原子操作

atomicAnd(int* address, int val);   // 位与atomicOr(int* address, int val);    // 位或atomicXor(int* address, int val);   // 位异或

3.原子操作与共享内存

可以操作共享内存，但是需要指定地址

__global__ voidsharedMemoryAtomic(int* globalData, int N) {    __shared__ int sharedCounter;    if (threadIdx.x == 0) {        sharedCounter = 0;    }    __syncthreads();    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        // 对共享内存执行原子操作        atomicAdd(&sharedCounter, 1);    }    __syncthreads();    // 将结果写回全局内存    if (threadIdx.x == 0) {        atomicAdd(&globalData[blockIdx.x], sharedCounter);    }}

4.性能优化建议

减少竞争：过多的线程竞争同一内存地址会严重影响性能

使用层次化方法：

线程块内使用共享内存进行部分归约

再使用原子操作更新全局内存

选择合适的数据类型：在满足需求的前提下，使用更小的数据类型

避免不必要的原子操作：如果不需要跨线程同步，不要使用原子操作

5.原子操作的局限性

性能开销：原子操作比普通内存操作慢得多

序列化：对同一内存地址的原子操作是串行执行的

不支持复杂操作：原子操作只支持基本算术和位运算

7.原子操作 vs 规约操作

cuda code

#include<iostream>#include<cuda_runtime.h>// 使用原子操作的朴素实现__global__ voidnaiveAtomicSum(int* data, int* result, int N){    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < N) {        atomicAdd(result, data[idx]);    }}// 使用共享内存优化的实现__global__ voidoptimizedAtomicSum(int* data, int* result, int N){    __shared__ int blockSum[256];    int tid = threadIdx.x;    int idx = blockIdx.x * blockDim.x + threadIdx.x;    blockSum[tid] = 0;    if (idx < N) {        blockSum[tid] = data[idx];    }    __syncthreads();    // 在共享内存中归约    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {        if (tid < stride) {            blockSum[tid] += blockSum[tid + stride];        }        __syncthreads();    }    // 只有线程0原子更新全局结果    if (tid == 0) {        atomicAdd(result, blockSum[0]);    }}intmain(){    const int N = 1000000;    int* h_data = new int[N];    int* d_data, *d_result;    // 初始化数据    for (int i = 0; i < N; i++) {        h_data[i] = 1;    }    // 分配设备内存    cudaMalloc(&d_data, N * sizeof(int));    cudaMalloc(&d_result, sizeof(int));    // 拷贝数据到设备    cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice);    // 测试朴素原子操作    int h_result = 0;    cudaMemset(d_result, 0, sizeof(int));    naiveAtomicSum<<<(N+255)/256, 256>>>(d_data, d_result, N);    cudaMemcpy(&h_result, d_result, sizeof(int), cudaMemcpyDeviceToHost);    std::cout << "Naive atomic sum: " << h_result << std::endl;    // 测试优化版本    h_result = 0;    cudaMemset(d_result, 0, sizeof(int));    optimizedAtomicSum<<<(N+255)/256, 256>>>(d_data, d_result, N);    cudaMemcpy(&h_result, d_result, sizeof(int), cudaMemcpyDeviceToHost);    std::cout << "Optimized atomic sum: " << h_result << std::endl;    // 清理    delete[] h_data;    cudaFree(d_data);    cudaFree(d_result);    return 0;}