首页 > 代码库 > CUDA: 原子操作

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

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

CUDA: 原子操作