共享内存实际上是可受用户控制的一级缓存。每个SM中的一级缓存与共享内存共享一个64KB的内存段在开普勒架构的设备中，根据应用程序的需要，每个线程块可以配置为16KB的一级缓存或共享内存。而在费米架构的设备中，可以根据喜好选择16KB或者48KB的一级缓存或者共享内存。早期费米架构中只有固定的16KB共享内存而没有一级缓存。共享内存的延迟极低，大约有1.5TB/s的带宽，远远高于全局内存的190GB/s，但是它的速度只有寄存器的十分之一。只有当数据重复利用，全局内存合并，或线程之间由数据共享时使用共享内存才更合适，否则，数据直接从全局内存加载到寄存器性能会更好。申请共享内存后，其内容在每一个用到的block被复制一遍，使得在每个block内，每一个thread都可以访问和操作这块内存，而无法访问其他block内的共享内存。这种机制就使得一个block之内的所有线程可以互相交流和合作。

  共享内存是基于存储体切换的架构（bank-switched architecture）。无论有多少个线程发起操作，每个存储体每个周期只执行一次操作。因此，如果线程束中的每个线程访问一个存储体，那么所有的线程操作都可以在一个周期内同时执行。此时无须顺序地访问，因为每个线程访问的存储体在共享内存中都是独立的，互不影响。此外，如果线程束中的所有线程同时访问相同地址的存储体时，就会想常量内存一样触发一个广播机制到线程束中的每一个线程。然而如果有其他的访问方式，存储体冲突将不同程度地得到解决。这意味着，线程访问共享内存需要排队等候，当一个线程访问时，其它线程都将阻塞。因此储存体应尽可能避免冲突。

                         

                             

          这里有必要对第四张图说明下。所有线程访问同一存储体，如果是读的话会触发广播机制，但是如果是写的话会冲突，这将导致对同一存储体顺序进行32次访问操作，这里注意箭头朝向，该图是参考博文https://segmentfault.com/a/1190000007533157，但是可能博主没考虑到读写的区别吧，如有问题，欢迎指正修改。

我们来看个向量点乘的例子吧。

#include "../common/book.h"

#define imin(a,b) (a<b?a:b)

const int N = 33 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


int main( void ) {
    float   *a, *b, c, *partial_c;
    float   *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the cpu side
    a = (float*)malloc( N*sizeof(float) );
    b = (float*)malloc( N*sizeof(float) );
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
                              N*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
                              N*sizeof(float) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // fill in the host memory with data
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    // copy the arrays 'a' and 'b' to the GPU
    HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float),
                              cudaMemcpyHostToDevice ) );
    HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float),
                              cudaMemcpyHostToDevice ) ); 

    dot<<<blocksPerGrid,threadsPerBlock>>>( dev_a, dev_b,
                                            dev_partial_c );

    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaMemcpyDeviceToHost ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    #define sum_squares(x)  (x*(x+1)*(2*x+1)/6)
    printf( "Does GPU value %.6g = %.6g?\n", c,
             2 * sum_squares( (float)(N - 1) ) );

    // free memory on the gpu side
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_partial_c ) );

    // free memory on the cpu side
    free( a );
    free( b );
    free( partial_c );
}

  主函数中的CUDA 内存申请，拷贝，释放我们就不具体介绍了，我们看核函数中的__shared__ float cache[threadsPerBlock] 这里的__shared__同样是CUDA中的关键字，和__global__一样，意思是申请一个共享内存，这里申请的大小为每个线程块的线程数目，这样每个线程都可以访问一个存储体，不会造成存储体冲突。并且会为每一个线程块分配相同大小的共享内存。这样我们就可以计算每个线程块的向量点乘之和，但是这里有个潜在的危险，如下图所示，如果一部分线程已经计算好了，另一部分还没，那么相加的结果肯定是错的。所以我们需要一个同步机制来保证每个线程块都计算好了，这里就用到了__syncthreads();

                                                                                         

                                             

总结：本次主要介绍了块内线程如何通过共享内存进行协作，除了线程级通信之外，我们之后还将学校线程块之间的通信。

参考：  [1]CUDA并行程序设计 —— GPU编程指南

      [2] Addison.Wesley.CUDA.By.Example.Jul.2010.ISBN.0131387685

      [3]https://segmentfault.com/a/1190000007533157