Linux系统下CUDA中的原子操作

今天在学习《CUDA并行程序设计》第四章关于统计直方图的例子时，在内核函数中调用了原子操作atomicAdd，以实现多个线程对竞争资源的使用。但是在编译时出现错误“ error: identifier "atomicAdd" is undefined”。该错误解决方案如下：

（1）确定你的GPU计算能力是否在1.1或者1.1以上？如果否，那你的硬件不支持，问题无法解决；如果是，看第（2）步

（2）在用nvcc编译时添加编译选项”-arch=sm_12"，即形如 ”nvcc -arch=sm_12 -o outFilename filename“这样的编译语句。另外如果你的GPU计算能力在2.0以上，也可以添加成“-arch=sm_20"编译选项。

经过上述步骤的检查，我的问题已经完美解决。

我程序的源代码如下：

__global__ void myhistogram_1(char * gpu_hist_data, int * gpu_bin_data)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	int idy = blockIdx.y * blockDim.y + threadIdx.y;
	int id  = idy * gridDim.x * blockDim.x + idx;
	char value;
	
	while(id < N)
	{
		value = gpu_hist_data[id];
		//gpu_bin_data[value] ++;
		atomicAdd(&gpu_bin_data[value], 1);
		id += blocksPerGrid * threadsPerBlock;
	}
	
}