关于CUDA共享存储器

Question

我使用的报告可用的共享存储器作为16384个字节

鉴于以下内核代码，其中T运行= INT（4个字节）

template <typename T> 
__global__ void foo(unsigned n, T *x) 
{ 
    unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; 
    __shared__ T sx[4096]; 
    x[idx] = 0; 
} 

我得到一个GTX9800奇怪的行为预期的结果，即最初非零的数组x将被零填充。

然而，添加一行代码，不会做任何事情：

template <typename T> 
__global__ void foo(unsigned n, T *x) 
{ 
    unsigned idx = blockIdx.x * blockDim.x + threadIdx.x; 
    __shared__ T sx[4096]; 
    sx[0] = 0; 
    x[idx] = 0; 
} 

现在X调用内核后不包含任何零可言！

但是，如果我将sx的大小更改为< = 4088，我会再次获得预期结果。

怎么回事？我很困惑。

编辑：

校正错字：改变16384 “KB” 到 “字节”

Answer 1

上计算能力1.x设备共享存储器的大小是16384个字节每 SM，不千字节。

此外，每个块将为内部目的消耗16个字节（存储块索引等），以及用于内核参数的附加存储。

因此，不幸的是，你不能在一个块中使用共享内存的全部16kb。

对于更高的计算能力，这些数据将存储在别处（常量内存和特殊寄存器），因此整个共享内存在那里可用。