CUDA： 原子操作

　　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]);
}  

通过降低内存竞争程度的策略来提高性能。