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);
共享内存
最新推荐文章于 2023-06-14 00:19:15 发布