CUDA的原子操作可以理解为对一个变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程,这个执行过程不能够再分解为更小的部分,在它执行过程中,不允许其他并行线程对该变量进行读取和写入的操作。基于这个机制,原子操作实现了对在多个线程间共享的变量的互斥保护,确保任何一次对变量的操作的结果的正确性。
**原子操作确保了在多个并行线程间共享的内存的读写保护,每次只能有一个线程对该变量进行读写操作,一个线程对该变量操作的时候,其他线程如果也要操作该变量,只能等待前一线程执行完成。**原子操作确保了安全,代价是牺牲了性能。
CUDA支持多种原子操作,常用的如下:
1、 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 处的32 位或64 位字old,计算(old + val),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64 位字。
2、 atomicSub()
int atomicSub(int* address, int val);
unsigned int atomicSub(unsigned int* address, unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算(old - val),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
3、 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);
读取位于全局或共享存储器中地址address 处的32 位或64 位字old,并将val 存储在存储器的同一地址中。这两项操作在一次原子事务中执行。该函数将返回old。只有全局存储器支持64 位字。
4、 atomicMin()
int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算old 和val 的最小值,并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
5、 atomicMax()
int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算old 和val 的最大值,并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
6、 atomicInc()
unsigned int atomicInc(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算 ((old >= val) ? 0 : (old+1)),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
7、 atomicDec()
unsigned int atomicDec(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算 (((old == 0) | (old > val)) ? val : (old-1)),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
8、 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);
读取位于全局或共享存储器中地址address 处的32 位或64 位字old,计算 (old == compare ? val : old),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old(比较并交换)。只有全局存储器支持64 位字。
9、 atomicAnd()
int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算 (old & val),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
10、 atomicOr()
int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算 (old | val),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
11、 atomicXor()
int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,unsigned int val);
读取位于全局或共享存储器中地址address 处的32 位字old,计算 (old ^ val),并将结果存储在存储器的同一地址中。这三项操作在一次原子事务中执行。该函数将返回old。
举个例子,定义1024个线程,求这1024个线程的ID之和,每个线程都会访问总和变量sum,如果不加原子操作,执行结果是错误并且是不确定的。
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#define SIZE 1024
__global__ void histo_kernel(int size, unsigned int *histo)
{
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < size)
{
//*histo+=i;
atomicAdd(histo, i);
}
}
int main(void)
{
int threadSum = 0;
//分配内存并拷贝初始数据
unsigned int *dev_histo;
cudaMalloc((void**)&dev_histo, sizeof(int));
cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);
// kernel launch - 2x the number of mps gave best timing
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
int blocks = prop.multiProcessorCount;
//确保线程数足够
histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);
//数据拷贝回CPU内存
cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);
printf("Threads SUM:%d\n", threadSum);
getchar();
cudaFree(dev_histo);
return 0;
}
使用原子操作正确的结果是523776,不使用原子操作的结果不确定,其中一次执行结果是711,显然是不对的。