__shared__ memory 的奇怪问题

[chinacodec]

一个 kernel, 我开了1024个 block, 256 个thread, GPU = 9300GM

__global__ void foo()
{
extern __shared__ float shared[];
}
这样写就可以成功,如果按照下面的写法, 程序就不执行foo,进去后马上退出foo程序
__global__ void foo()
{
__shared__ float shared[2048];
}
而我实际使用的 shared memory 大小就是2048, 请问有没有人碰到过类似问题

[qin_y_j]

sharemem 最大8KB；
参数也会占用一小部分空间；
2048 * sizeof（float） = 8192B太大了；
试试2000一下的；

[lysj9]

更正下，shared memory最多可以到16KB
你说的实际使用的shared memory大小是2048，是2048个float还是2048B？，前者的话要除4才对哎……
用extern的情况下，外部指定shared memory大小，一般就指定类似2048B这样吧，比如kernel<<<gridDim,blockDim,2048>>>();

[lysj9]

错了，是2048B才要/4

[chinacodec]

__shared__ float shared[2048];

而我实际使用的 shared memory 大小就是2048个float，即8K

按照nvidia官方说有16k，应该是足够。这个是其一
我用 extern __shared__ float shared[]来分配的话，也访问到了 shared[2047]这个单元，所以说共享内存是应该足够多的，但是我手工指定就无法正常工作了

[lysj9]

后者调用kernel时候没指定shared memory大小吧
不解……等高人回复吧

[naughty_iori]

看看出错时的ERROR　CODE.
第一种情况要看你调用KERNEL的时候在外面指定的SHARE　MEM有多大。如果大于8192的话，我觉得有可能是KERNEL里有问题。比如内存访问越界，因为第一种里的内存较大，所以没问题。第二种就是固定那么大，越界了就是出错，这时ERROR　CODE应该是CUDA_ERROR_LAUNCH_FAILED。

如果第一种指定的是8192的话，我也不知道是什么问题了。如果SHARE　MEM不够用的话ERROR　CODE应该是CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES.