静态分配
加上前缀 shared
__shared__ int _ss[1024];
- 1
动态分配
当我们在编程时,不清楚shared memory 数组开多大,就要用到动态分配。
分为两部分:
1, 声明
extern __shared__ int _s[];
- 1
2, 在调用kernel 时加上数组的大小。
xxx_kernel<<<grid, block, sharedMemSize>>>();
- 1
内存分布
下面通过一个例子,说明同时使用静态和动态分配时, 内存分配情况。
kernel 代码:
__global__ void sharedMemTest()
{
extern __shared__ int _s[];
__shared__ int _ss[1024];
if (threadIdx.x==0)
printf("blockIdx.x is %d s is at %x, ss is at %x\n", blockIdx.x, _s, _ss);
}
调用kernel代码:
{
dim3 block(32);
dim3 grid(32);
sharedMemTest << <grid, block, 4*1024>> >();
cudaDeviceSynchronize();
}
输出结果如下:
blockIdx.x is 27 s is at 1001000, ss is at 1000000
blockIdx.x is 6 s is at 1001000, ss is at 1000000
blockIdx.x is 9 s is at 1001000, ss is at 1000000
blockIdx.x is 18 s is at 1001000, ss is at 1000000
…
blockIdx.x is 30 s is at 1001000, ss is at 1000000
blockIdx.x is 10 s is at 1001000, ss is at 1000000
可以看出以下几点:
1, 每个block 都有自己独立的shared memory地址空间。
2, 静态开辟的空间总是从地址1000000开始。
3, 动态开辟空间是在静态空间之后的。
如果将动态开辟地址大小设置太大,导致整个block 使用的shared memory 空间超过maxSharedMemoryPerBlock,会导致kernel 不执行。例如将调用代码改成下面:
{
dim3 block(32);
dim3 grid(32);
sharedMemTest << <grid, block, 48*1024>> >();
cudaDeviceSynchronize();
}
由于我的显卡中maxSharedMemoryPerBlock = 48KB,动态空间+静态 = 49KB所以程序并没有输出。
原文链接:https://blog.csdn.net/galaxy_wolf/article/details/50738200