cuda编程笔记（5）--原子操作
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) {// 释放锁：把锁设为 0atomicExch(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和histochar* 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;
}