码迷,mamicode.com
首页 > 其他好文 > 详细

CUDA: 原子操作

时间:2017-04-21 16:44:07      阅读:254      评论:0      收藏:0      [点我收藏+]

标签:通过   支持   atomic   lob   __sync   操作   性能   atom   全局   

  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: 原子操作

标签:通过   支持   atomic   lob   __sync   操作   性能   atom   全局   

原文地址:http://www.cnblogs.com/programmer-wfq/p/6744115.html

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!