一、知识梳理
在实际的开发过程中,因为CUDA核函数是在每个线程上并行执行的,在对其中的变量进行写操作的时候,很容易发生BUG,所以CUDA推出了一个名为“原子操作”的机制,这个机制类似Linux的“原子操作”,通俗来讲就是为参与运算的变量加了一个锁,线程对变量进行操作需要有上锁、解锁的操作,当有线程持有锁时,其他线程如果要对变量进行操作就需要等待持有锁线程解锁。
![](https://i-blog.csdnimg.cn/blog_migrate/0f4dbd9f14f3a596ac2d12a872f723dd.png)
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;
}
![](https://i-blog.csdnimg.cn/blog_migrate/c78617d26356062a7079f2397a1c022d.png)
使用原子操作
![](https://i-blog.csdnimg.cn/blog_migrate/4b4aff333a97497f165bd15dd0838cc9.png)
不使用原子操作
可以看到,不使用原子操作会使结果产生严重错误。
三、总结
因为CUDA核函数是并行执行的,所以要格外注意对变量的保护,不然就是失之毫厘,谬之千里。