CUDA原子操作【参加CUDA线上训练营】

目录

CUDA原子操作的概念

实例：向量元素求和 - GPU方法

原子操作常用函数

代码示例

---

CUDA原子操作的概念

起因：多个线程访问同一个资源，造成的冲突。

 - CUDA的原子操作可以理解为对一个Globalmemory或Sharedmemory中变量进行“读取-修改-写入”这三个操作的一个最小单位的执行过程（含义），在它执行过程中，不允许其他并行线程对该变量进行读取和写入的操作。

 - 基于这个机制，原子操作实现了对在多个线程间共享的变量的互斥保护（意义/作用），确保任何一次对变量的操作的结果的正确性。当前资源被一个线程占用，那其他线程需要此资源就要进行等待。

实例：向量元素求和 - GPU方法

步骤：

1、申请N个线程

2、通过公式计算得到当前线程在所有线程中的index

3、每个线程读取一个数据，并放到所在block中的shared memory中

4、利用__synthreads()同步，等待所有线程执行完毕

5、每个线程读取所在block中shared memory中的数据，每次读取两个做加法，同步直到所有线程都做完，并将结果写到它所对应的shared memory位置中

6、迭代进行第5步，每次迭代步数减少一半，直到只剩一个线程。这样所有的shared memory当中的数值都累加完毕

7、每个线程块的shared memory中第0号位置，就保存了该线程块中所有数据的总和

8、将每个block的shared memory中第一个值放到输出向量里面

9A、再执行第二次核函数，只需要一个block，并且将所有数据放进shared memory中，重复之前的步骤，最后输出结果。

 

9B、不需要进行第二次核函数，直接在第一次核函数中的最后，利用原子操作，直接进行累加。

 

原子操作常用函数

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);
float atomicAdd(float* address, float val);
double atomicAdd(double* address, double val);
__half2 atomicAdd(__half2 *address, __half2 val);
__half atomicAdd(__half *address, __half val);
__nv_bfloat162 atomicAdd(__nv_bfloat162 *address, __nv_bfloat162 val);
__nv_bfloat16 atomicAdd(__nv_bfloat16 *address, __nv_bfloat16 val);
读取位于全局或共享内存中地址 address 的 16 位、32 位或 64 位字 old，计算 (old + val)，并将结果存储回同一地址的内存中。这三个操作在一个原子事务中执行。该函数返回old。

atomicAdd() 的 32 位浮点版本仅受计算能力 2.x 及更高版本的设备支持。

atomicAdd() 的 64 位浮点版本仅受计算能力 6.x 及更高版本的设备支持。

atomicAdd() 的 32 位 __half2 浮点版本仅受计算能力 6.x 及更高版本的设备支持。 __half2 或 __nv_bfloat162 加法操作的原子性分别保证两个 __half 或 __nv_bfloat16 元素中的每一个；不保证整个 __half2 或 __nv_bfloat162 作为单个 32 位访问是原子的。

atomicAdd() 的 16 位 __half 浮点版本仅受计算能力 7.x 及更高版本的设备支持。

atomicAdd() 的 16 位 __nv_bfloat16 浮点版本仅受计算能力 8.x 及更高版本的设备支持。

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。

4. atomicMin()
int atomicMin(int* address, int val);
unsigned int atomicMin(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMin(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMin(long long int* address,
                                long long int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old，计算 old 和 val 的最小值，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。

atomicMin() 的 64 位版本仅受计算能力 3.5 及更高版本的设备支持。

5. atomicMax()
int atomicMax(int* address, int val);
unsigned int atomicMax(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicMax(unsigned long long int* address,
                                 unsigned long long int val);
long long int atomicMax(long long int* address,
                                 long long int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old，计算 old 和 val 的最大值，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。

atomicMax() 的 64 位版本仅受计算能力 3.5 及更高版本的设备支持。

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);
unsigned short int atomicCAS(unsigned short int *address, 
                             unsigned short int compare, 
                             unsigned short int val);
读取位于全局或共享内存中地址address的 16 位、32 位或 64 位字 old，计算 (old == compare ? val : old) ，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old（Compare And Swap）。

Bitwise Functions
9. atomicAnd()
int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
                                 unsigned long long int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old，计算 (old & val)，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。

atomicAnd() 的 64 位版本仅受计算能力 3.5 及更高版本的设备支持。

10. atomicOr()
int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
                      unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
                                unsigned long long int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old，计算 (old | val)，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。

atomicOr() 的 64 位版本仅受计算能力 3.5 及更高版本的设备支持。

11. atomicXor()
int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
                                 unsigned long long int val);
读取位于全局或共享内存中地址address的 32 位或 64 位字 old，计算 (old ^ val)，并将结果存储回同一地址的内存中。 这三个操作在一个原子事务中执行。 该函数返回old。

atomicXor() 的 64 位版本仅受计算能力 3.5 及更高版本的设备支持。

代码示例

#include<stdio.h>
#include<stdint.h>
#include<time.h>     //for time()
#include<stdlib.h>   //for srand()/rand()
#include<sys/time.h> //for gettimeofday()/struct timeval
#include"error.cuh"

#define N 10000000
#define BLOCK_SIZE 256
#define BLOCKS 32 


__managed__ int source[N];               //input data
__managed__ int final_result[1] = {0};   //scalar output

__global__ void _sum_gpu(int *input, int count, int *output)
{
   __shared__ int sum_per_block[BLOCK_SIZE];
   int temp = 0;
   int idx = blockDim.x * blockIdx.x + threadIdx.x;
   for(;idx < count;idx += gridDim.x * blockDim.x)//数据大于线程数量，用for循环先将每个线程每次需要处理的数据先相加
   {
       temp += input[idx];//将每个线程每次需要处理的数据相加
   }
   sum_per_block[threadIdx.x] = temp;  //将和传给shared memory中，即每个线程的和数据
    __syncthreads();//等待所有线程完毕

   for (int length = BLOCK_SIZE / 2; length >= 1; length /= 2)//每个block中线程按length步长相加
   {
       int tmp_sum = 0;
       if(threadIdx.x < length)//小于步长的线程进行求和运算
       {
           tmp_sum = sum_per_block[threadIdx.x] + sum_per_block[threadIdx.x + length];
       }
       __syncthreads();//等待所有线程求和完毕
       
       if (threadIdx.x < length)
        {
            sum_per_block[threadIdx.x] = tmp_sum;//将临时储存的结果替换shared memory中0号的数据
        }
        __syncthreads();
   }
   
   if (blockDim.x * blockIdx.x < count) //防止block和block_size设置错误
    {
        //the final reduction performed by atomicAdd()
        if (threadIdx.x == 0) atomicAdd(output, sum_per_block[0]);//将每个block中0号的shared memory累加
    }
}

int _sum_cpu(int *ptr, int count)
{
    int sum = 0;
    for (int i = 0; i < count; i++)
    {
        sum += ptr[i];
    }
    return sum;
}

void _init(int *ptr, int count)
{
    uint32_t seed = (uint32_t)time(NULL); //make huan happy
    srand(seed);  //reseeding the random generator

    //filling the buffer with random data
    for (int i = 0; i < count; i++) ptr[i] = rand();
}

double get_time()
{
    struct timeval tv;
    gettimeofday(&tv, NULL);
    return ((double)tv.tv_usec * 0.000001 + tv.tv_sec);
}

int main()
{
    //**********************************
    fprintf(stderr, "filling the buffer with %d elements...\n", N);
    _init(source, N);

    //**********************************
    //Now we are going to kick start your kernel.
    cudaDeviceSynchronize(); //steady! ready! go!
    
    fprintf(stderr, "Running on GPU...\n");
    
double t0 = get_time();
    _sum_gpu<<<BLOCKS, BLOCK_SIZE>>>(source, N, final_result);
    CHECK(cudaGetLastError());  //checking for launch failures
    CHECK(cudaDeviceSynchronize()); //checking for run-time failurs
double t1 = get_time();

    int A = final_result[0];
    fprintf(stderr, "GPU sum: %u\n", A);


    //**********************************
    //Now we are going to exercise your CPU...
    fprintf(stderr, "Running on CPU...\n");

double t2 = get_time();
    int B = _sum_cpu(source, N);
double t3 = get_time();
    fprintf(stderr, "CPU sum: %u\n", B);

    //******The last judgement**********
    if (A == B)
    {
        fprintf(stderr, "Test Passed!\n");
    }
    else
    {
        fprintf(stderr, "Test failed!\n");
	exit(-1);
    }
    
    //****and some timing details*******
    fprintf(stderr, "GPU time %.3f ms\n", (t1 - t0) * 1000.0);
    fprintf(stderr, "CPU time %.3f ms\n", (t3 - t2) * 1000.0);

    return 0;
}	
	

运行结果：