ref docs.nvidia.com.shared-memory

1. 共享内存在哪里

下图为Turing架构GPU的SharedMemory位置。是片上内存中最大的一片。它的存在以SM为单位。

2. 共享内存有什么特点

共享内存的位置不同于全局内存，在片上拥有更快的访问速度，访问速度比未缓存的全局内存快100倍。
共享内存的硬件设计被划分为可以同时访问的同等大小的内存模块。因此，可以同时处理跨越 n 个不同内存组的任何内存负载或n个地址的存储，从而产生比单个存储体的带宽高n倍的有效带宽。

正确的运用这些特点能帮助更高效的执行特定方式的访存。

3. 应用于线程间并发访存

3.1. 并发访问的特性

根据共享内存的访问属性。当不同线程的内存访问被合理分配到不同内存块(Bank)上，有效带宽达到最大。反之，当不同线程访问冲突的内存块，访问被序列化，有效带宽会减少一个系数。
此处的一个例外是，当 warp 中的多个线程对同一内存块的同一个位置访问会触发广播。广播会将多个访问并发。

3.2. 内存块(Bank)的大小

在计算能力3.x以上的设备（目前GeForce GTX 1060 的计算能力是6.1）具有可配置的Bank大小。通过设置 cudaSharedMemBankSizeEightByte / cudaSharedMemBankSizeFourByte 控制Bank大小。其中 cudaSharedMemBankSizeEightByte 有助于避免访问双精度数据时的Bank时冲突。
这里我们也可以看到，当warp发射线程数为32、bank数量大于32时，并行访问不同寻址的共享内存不会发生冲突。

3.3. 访存的线程间同步

竞争访问SharedMemory会导致未定义行为。
例如，warp表面上并行发射线程，实际上分批发射线程。故，线程A与线程B并发读取全局内存到共享内存后后，线程A访问B共享内存有可能发生冲突。
这里CUDA引入 __syncthreads() 同步原语使并行发射的函数在这里等待直至并行发射的函数全部执行到这里。

3.4. 静态 / 动态共享内存申请和使用

这里的静态、动态指的是共享内存的大小在编译器明确或运行期明确。他们的使用区别如下

静态共享内存

// 发射cuda kernel
staticReverse<<<1,n>>>(ptr, size);

// kernel
__global__ void staticReverse(int *d, int n) {
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

这里使用关键字 __shared__

动态共享内存

// 发射 cuda kernel
dynamicReverse<<<1,n,n*sizeof(int)>>>(ptr, n);

// kernel
__global__ void dynamicReverse(int *d, int n) {
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

这里使用关键字 extern __shared__ 。需要注意，当kernel函数中需要多块共享内存数组时，只能调用一次 extern __shared__ 并手动分割此内存块用于接下来的计算。

4. 共享内存与全局内存的数据传输

CUDA 11.0 引入了 异步复制功能。通过异步拷贝，能绕过RF寄存器和L1Cache传输数据。示例代码如下

//pipeline pipe;
for (size_t i = 0; i < copy_count; ++i) {
  __pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                          &global[blockDim.x * i + threadIdx.x], sizeof(T));
}

4.1.1 【NVIDIA-GPU-CUDA】片上的并发访问存储 —— Shared Memory