共享内存

最新推荐文章于 2023-06-14 00:19:15 发布

scut许sir

最新推荐文章于 2023-06-14 00:19:15 发布

阅读量463

点赞数

分类专栏： CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/u010053524/article/details/50747791

版权

9 篇文章 2 订阅

订阅专栏

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);

关注

专栏目录