cuda共享内存的作用和功能验证

在cuda中定义共享内存需要使用关键字__shared__
在一个核函数里面声明一个共享内存，其它的线程也可以访问它。
共享内存是不能被初始化的，只能先定义再赋值。
下面是验证共享内存功能的一个例子。

#include<iostream>
#define MEM_SIZE 10

__global__ void shareVisit(int* gpuMem)
{
    //定义共享变量
    __shared__ int shareValue;
    //当前位置的线程编号
    int thisId=threadIdx.x;
    //判断是不是第1个线程,如果是第1个线程，对线程做初始化
    if(thisId==0) shareValue=MEM_SIZE+1;
    //同步线程，保证都运行到了这个位置再往后走
    __syncthreads();
    //把共享内存里面的数据写入到当前位置
    gpuMem[thisId]=shareValue;
    
}

int main()
{
    //新建一个gpu的内存
    int* gpuMem;
    cudaMalloc((void**)&gpuMem,sizeof(int)*MEM_SIZE);
    //调用共享内存的测试核函数
    shareVisit<<<1,MEM_SIZE>>>(gpuMem);
    //新建一个cpu的内存
    int cpuMem[MEM_SIZE];
    //把gpu里面的数据复制给cpu
    cudaMemcpy(cpuMem,gpuMem,sizeof(int)*MEM_SIZE,cudaMemcpyDeviceToHost);
    //遍历cpu里面的数据并且依次显示出来
    for(int i=0;i<MEM_SIZE;++i)
    {
        std::cout<<cpuMem[i]<<std::endl;
    }
    return 0;
}

最后会打印出10个"11"
另外需要注意，共享内存仅仅是同一个block内的共享，不同的block之间的共享内存是不会相互影响的，下面的代码可以在一定程度上说明这个问题。

#include<iostream>

__global__ void shareTest(int* gpuData)
{
    //获取当前位置的线程编号
    int thrId=threadIdx.x;
    //新建共享内存
    __shared__ int shareData;
    //判断是不是第1个线程
    if(thrId==0)
    {
        //令第1个线程等于block的标号
        shareData=blockIdx.x;
    }
    //做同步，确保每个线程都运行到了这里
    __syncthreads();
    //在对应的位置写上线程的共享内存
    gpuData[blockIdx.x*blockDim.x+thrId]=shareData;
}

int main()
{
    int* gpuData;
    cudaMalloc((void**)&gpuData,sizeof(int)*20);
    //在gpu数据里面写值
    shareTest<<<2,10>>>(gpuData);
    //把数据复制到cpu里面
    int cpuData[20];
    cudaMemcpy(cpuData,gpuData,sizeof(int)*20,cudaMemcpyDeviceToHost);
    //依次输出每个位置的数据
    for(int i=0;i<2;++i)
    {
        for(int j=0;j<10;++j) std::cout<<cpuData[i*10+j]<<" ";
        std::cout<<std::endl;
    }
}

最后会输出

0 0 0 0 0 0 0 0 0 0
1 1 1 1 1 1 1 1 1 1