CUDA学习之路2 共享内存的使用

最新推荐文章于 2025-02-08 11:55:32 发布

xiaobaicai666

最新推荐文章于 2025-02-08 11:55:32 发布

阅读量1.5k

点赞数 2

CC 4.0 BY-SA版权

分类专栏： CUDA 文章标签： CUDA 共享内存

本文链接：https://blog.youkuaiyun.com/xiaobaicai666/article/details/85099077

CUDA 专栏收录该内容

2 篇文章

订阅专栏

本文深入探讨了CUDA编程中共享内存的正确使用方法，强调了线程同步的重要性，通过一个归约求和的实例，展示了忽视__syncthread()函数导致的数据计算错误，提醒开发者在使用共享内存时务必加入线程同步。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

这两天在写一个对任意长度数组进行扫描的cuda代码，由于对共享内存的理解不清楚，浪费了好多时间，写篇博客记录一下，避免下次再犯。

正文：

当我们使用共享内存时，要谨记，但凡涉及到共享内存的操作，一定要记得使用线程同步__syncthread();否则必然会造成数据的计算出错。因为共享内存在块内对所有线程都是可见的，是大家一起使用的。如果你在每次使用之后不同步一下，那么当运算的快的线程运算完成后，就可能会读取还未完成运算的数据进行下一次的运算了。

下面这段代码是一个归约求和的例子，如果我们对[0,1,2,3,4,5,6,7]进行求和，如果共享内存使用后没有加那两个__syncthread();那么就可能发生线程冲突。例如第一次归约后我们应该得到[4,6,8,10,4,5,6,7]然后再进行下一步的归约，但是若没有加线程同步，则可能tid == 0的线程先运算完第一轮，得到0+4=4；然后紧接着开始下一轮，由于其他线程还在计算，还未把结果放回原来的位置，这个时候索引为 2 的位置上数字还是2，则下一次运算的就是4+2=6而不是4+8=12了。

__global__ void reduction(int* d_in,int arr_length){

    int tid = threadIdx.x;
    int idx = blockDim.x*blockIdx.x + threadIdx.x;
    __shared__ int sdata[512];
    sdata[tid] = (idx < arr_length) ? d_in[idx]:0;
    __syncthread();//这里必须加同步，否则下面的归约求和会乱
    for(int s = blockjDim.x / 2; s>0; s>>=1){
        if(tid < s){
            sdata[tid] += sdata[tid + s];
        }
        __syncthread();
    }
}

下次再忘记加__syncthread();我就是狗！！！