共享内存

• static local shared memory定义在kernel函数中，只在该kernel中可见

  __global__ void kernelFunc() {
      __shared__ Type data;
      ...
  }
  

• static global shared memory定义在文件中的任何kernel外，对所有kernel可见

  __shared__ Type data;
  __global__ void kernelFunc() {
      ...
  }
  

• dynamci shared memory要在变量声明前添加extern关键字，动态共享内存必须是一维的，可以定义为局部或全局的。一个kernel只能有一段动态共享内存，其大小在调用kernek时指定

  __global__ void kernelFunc() {
      extern __shared__ Type data[];
      ...
  }
  int main() {
      kernelFunc <<<gridDim,blockDim,sharedMemArrLength>>> ();
  }
  

• bank conflict指的是同一线程束内多个线程访问同一个bank对应的不同内存地址，通俗的讲，就是访问了不同行的同一列地址。bank width决定了内存操作的粒度，开普勒架构中，可以设置为4-byte或8-byte，较大的bank width会提高bank conflict的几率，但也能得到较高的带宽

  cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);
  cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);
  

• L1 cache和shared memory共享同一块on-chip memory，可以手动配置某个设备上所有或者单个SM上L1 cache和shared memory的大小

  cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);
  cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheca cheConfig);