参加【Nvidia CUDA线上训练营2023春】笔记——使用原子操作保护变量

一、知识梳理

在实际的开发过程中，因为CUDA核函数是在每个线程上并行执行的，在对其中的变量进行写操作的时候，很容易发生BUG，所以CUDA推出了一个名为“原子操作”的机制，这个机制类似Linux的“原子操作”，通俗来讲就是为参与运算的变量加了一个锁，线程对变量进行操作需要有上锁、解锁的操作，当有线程持有锁时，其他线程如果要对变量进行操作就需要等待持有锁线程解锁。

i为多少

二、实验示例

CUDA为我们提供了一系列的原子操作函数，如：

//加法       atomicAdd(&value，num)  value = valude + num
//减法       atomicSub(&value，num)  value = valude - num
//赋值       atomicExch(&value，num)  value = valude + num
//求最大     atomicMax(&value，num)  value = max(value,num)
//求最小     atomicMin(&value，num)  value = main(value,num)
//向上计数    atomicLnc(&value，num)  value = (value <= num) ? valude++ : 0;
//向下计数    atomicDec(&value，num)   value = (value > num)  ? value-- : 0;
//计较并交换  atomicCAS(&value，num)  value = valude & num
//与运算      atomicAnd(&value，num)  value = valude | num
//或运算      atomicOr(&value，num)  value = valude | num
//异或运算    atomicXor(&value，num)  value = valude ^ num

我们以加法为例，进行实验:

#include<stdio.h>
#include<math.h>
#include "error.cuh"

#define ARRAY_MAX 1000000

__global__ void array_add_d(int *d_a,int *d_out)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    
    atomicAdd(d_out,d_a[index]);
    //d_out[0] += d_a[index];

}

void array_add_h(int *h_a,int *h_out_h)
{
    int i;
    for(i = 0; i < ARRAY_MAX; i++)
    {
        h_out_h[0] += h_a[i];
    }
}


int main(int argc, char const *argv[])
{
    int *h_a,*h_out,*h_out_h;
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*ARRAY_MAX));
    CHECK(cudaMallocHost((void **) &h_out, sizeof(int)));
    CHECK(cudaMallocHost((void **) &h_out_h, sizeof(int)));
    
    for(int i = 0; i < 1000000; i++)
    {
        h_a[i] = 1;
    }
    *h_out = {0};
    *h_out_h = {0};
    h_a[3245] = 10;
    h_a[4556] = 0;
    
    int *d_a,*d_out;

    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*ARRAY_MAX));
    CHECK(cudaMalloc((void **) &d_out, sizeof(int)));
    
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*ARRAY_MAX, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_out, h_out, sizeof(int), cudaMemcpyHostToDevice));
    
    array_add_d<<<1000, 1000>>>(d_a, d_out);
    array_add_h(h_a,h_out_h);
    
    CHECK(cudaMemcpy(h_out, d_out, sizeof(int), cudaMemcpyDeviceToHost));
    
    if(h_out[0] == h_out_h[0])
    {
        printf("Count Correct!\n");
    }else
    {
        printf("Count Failed-_-\n");
    }
    
    printf("DATA_OUT is %d\n",h_out[0]);
    
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_out));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_out));
    
    
    
    return 0;
}

使用原子操作

不使用原子操作

可以看到，不使用原子操作会使结果产生严重错误。

三、总结

因为CUDA核函数是并行执行的，所以要格外注意对变量的保护，不然就是失之毫厘，谬之千里。