原子操作CUDA中的原子操作

简介: 原子操作CUDA中的原子操作

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,如果不加原子操作,执行结果是错误并且是不确定的。

[cpp]  view plain  copy

print?

  1. #include <stdio.h>      
  2. #include <stdlib.h>    
  3. #include <cuda_runtime.h>    
  4.  
  5. #define SIZE 1024  
  6.  
  7. __global__ void histo_kernel(int size, unsigned int *histo)  
  8. {  
  9.    int i = threadIdx.x + blockIdx.x * blockDim.x;  
  10.    if (i < size)  
  11.    {  
  12.        //*histo+=i;  
  13.        atomicAdd(histo, i);  
  14.    }  
  15. }  
  16.  
  17. int main(void)  
  18. {  
  19.    int threadSum = 0;  
  20.  
  21.    //分配内存并拷贝初始数据  
  22.    unsigned int *dev_histo;  
  23.  
  24.    cudaMalloc((void**)&dev_histo, sizeof(int));  
  25.    cudaMemcpy(dev_histo, &threadSum, sizeof(int), cudaMemcpyHostToDevice);  
  26.  
  27.    // kernel launch - 2x the number of mps gave best timing    
  28.    cudaDeviceProp  prop;  
  29.    cudaGetDeviceProperties(&prop, 0);  
  30.  
  31.    int blocks = prop.multiProcessorCount;  
  32.    //确保线程数足够  
  33.    histo_kernel << <blocks * 2, (SIZE + 2 * blocks - 1) / blocks / 2 >> > (SIZE, dev_histo);  
  34.  
  35.    //数据拷贝回CPU内存  
  36.    cudaMemcpy(&threadSum, dev_histo, sizeof(int), cudaMemcpyDeviceToHost);  
  37.    printf("Threads SUM:%d\n", threadSum);  
  38.    getchar();  
  39.    cudaFree(dev_histo);  
  40.    return 0;  
  41. }  

使用原子操作正确的结果是523776,不使用原子操作的结果不确定,其中一次执行结果是711,显然是不对的。

目录
相关文章
|
5月前
【原子操作】顺序操作
【原子操作】顺序操作
|
11月前
|
Cloud Native Go C语言
C 语言的 互斥锁、自旋锁、原子操作
C 语言的 互斥锁、自旋锁、原子操作
|
2月前
|
存储 算法
什么是原子操作?
【8月更文挑战第24天】
45 0
|
10月前
|
API 调度 C语言
互斥锁,自旋锁,原子操作的原理,区别和实现
v互斥锁,自旋锁,原子操作的原理,区别和实现
115 0
多线程的原子操作
多线程的原子操作
62 0
|
5月前
|
Linux C语言
linux c 多线程 互斥锁、自旋锁、原子操作的分析与使用
生活中,我们常常会在12306或者其他购票软件上买票,特别是春节期间或者国庆长假的时候,总会出现抢票的现象,最后总会有人买不到票而埋怨这埋怨那,其实这还好,至少不会跑去现场或者网上去找客服理论,如果出现了付款,但是却没买到票的现象,那才是真的会出现很多问题,将这里的票引入到多线程中,票就被称为临界资源。
74 0
|
5月前
|
Linux
【Linux C 几种锁的性能对比】 1.读写锁 2.互斥锁 3.自旋锁 4.信号量 5.rcu
【Linux C 几种锁的性能对比】 1.读写锁 2.互斥锁 3.自旋锁 4.信号量 5.rcu
|
5月前
|
安全 Linux 调度
Linux C/C++ 开发(学习笔记四):多线程并发锁:互斥锁、自旋锁、原子操作、CAS
Linux C/C++ 开发(学习笔记四):多线程并发锁:互斥锁、自旋锁、原子操作、CAS
80 0
|
5月前
|
编译器 C++
vs2017下原子操作和互斥锁性能比较
vs2017下原子操作和互斥锁性能比较
|
11月前
互斥锁、自旋锁、原子操作
互斥锁、自旋锁、原子操作