当前位置：首页 > news >正文

cuda编程笔记（5）--原子操作

news 2025/6/28 10:15:32

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;
}

查看全文

http://www.lqws.cn/news/544501.html