1.1以上计算功能集支持全局内存上的原子操作， 1.2以上支持共享内存上的原子操作。
atomicAdd(add,y)将生成一个原子的操作序列，这个操作序列包括读取地址addr处的值，将y增加到这个值，以及将结果保存回地址addr。
一个统计字符出现频率的直方图GPU内核函数：
__global__ void histo_kernel(unsigned char* buffer, long size, unsigned int* histo){ __shared__ unsigned int temp[256]; tmp[threadIdx.x] = 0; __syncThreads(); int i = threadIdx.x + blockIdx.x * blockDim.x; int offset = blockDim.x * gridDim.x; while(i<size){ atomicAdd( &temp[buffer[i]], 1); i += offset; } __syncthreads(); atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x]); }
通过降低内存竞争程度的策略来提高性能。
原文：http://www.cnblogs.com/programmer-wfq/p/6744115.html