CUDA共享内存详解

为什么需要共享内存？

共享内存的访问速度比访问全局速度快的多，因此对于多次访问全局内存的程序，特别是需要多次将全局内存的运算结果缓存到全局内存的运算，先将临时结果缓存到共享内存再做计算，会提高运算速度。

1、例如累积归并求和，将一串数字从头加到尾，归并求和算法如下：

如果使用全局内存，每做一次减半运算，都要访问一半的全局内存，总要访问的全局内存次数为 N + N/2 + N/4 +… + 1 次。因此减少全局内存的访问次数可以提高运算速度。
如果提高呢？那就是第一次先将全局内存复制到共享内存中，这时总需要访问N次全局内存，每个线程只需要访问一次。

2、再如直方图统计：
假设直方图统计的bin个数为256，假设统计N 个数，如果使用全局内存，需要将N个全局内存的数累加到256个全局内存中，那么也就是需要访问了2N次全局内存。如果将256 个bin的内存改为共享内存，那么只需要访问N次的全局内存，运算速度得到提升。

什么是共享内存

对于CUDA，一个grid 有多个Block 块，一个block块多个线程。
共享内存只是 Block 块内的线程共享，不同Block块之间的共享内存是不会共享的。
定义在核函数内，加前关键字为__shared__，例如：
shared int sharedata[thread_perblock];

块内同步

在做共享内存时，一般需要使用线程块内同步__syncthreads()，只是让相同block块中的线程都要达到这个位置后才可往下执行。

如何实现

为了使用共享内存，一般定义共享内存的大小为Block 块中的线程个数。这样每个线程都对应到响应的共享内存中。使用线程id 进行对应。
int thread_id = threadIdx.x;
对于不同块中的共享内存，之间不存在共享，因此全局映射到共享内存方式如下：
int tid = threadIdx.x + blockDim.x * blockIdx.x;
sharedata[thread_id] = data[tid];
tid 处于哪个Block 中即映射到哪个Block块中的共享内存。

做完映射和对应后即可进行运算，块内运算以块为基本单元，使用id为块内线程id:
例如规约求和：

 int mid = thread_perblock / 2;  //规约求和， thread_perblock 必须为 2 的 K 次方
    while(mid != 0){
        if(thread_id < mid){
            sharedata[thread_id] += sharedata[thread_id + mid];
        }
        __syncthreads();
        mid /= 2;
    }

每做完一轮块内运算进行一次同步，等待块内所有线程执行完。while里即为循环次数。
等所有块执行完后将共享内存的结果再赋值到全局内存，完成整个过程。

 out[blockIdx.x] = sharedata[0]; // 不同block中的shareddata 互不干扰

代码

1、归并求和：

#define thread_perblock 8
 __global__ void reduce_add(int *data, int *out, int N){
    __shared__ int sharedata[thread_perblock];
    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    int thread_id = threadIdx.x;
    printf("tid = %d, %d\n", tid, data[tid]);
    while (tid < N) // 数据预处理到共享内存，当数组大于线程数量，需要一个线程处理多个数据，跨度(blockDim.x * gridDim.x)
    {
        sharedata[thread_id] += data[tid];
        tid += blockDim.x * gridDim.x;
    }
    __syncthreads();// 块内同步

    int mid = thread_perblock / 2;  //规约求和， thread_perblock 必须为 2 的 K 次方
    while(mid != 0){
        if(thread_id < mid){
            sharedata[thread_id] += sharedata[thread_id + mid];
            printf("blockIdx.x = %d, sharedata[%d] = %d , mid = %d\n " ,blockIdx.x, thread_id, sharedata[thread_id], mid);
        }
        __syncthreads();
        mid /= 2;
    }
    out[blockIdx.x] = sharedata[0]; // 不同block中的shareddata 互不干扰
    printf("blockIdx.x = %d, sharedata[0] = %d \n" , blockIdx.x,  sharedata[0]);
 }

2、直方图统计：

__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo) {
    __shared__ unsigned int temp[256];
    temp[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]);
}